Updating changes needed to propogate the correct resolution into gst…

…Camera. Support for YUV

FindSDL2TTF.cmake

@@ -80,7 +80,7 @@ SET(SDL2TTF_SEARCH_PATHS
 FIND_PATH(SDL2TTF_INCLUDE_DIR SDL_ttf.h
 	HINTS
 	$ENV{SDL2TTFDIR}
-	PATH_SUFFIXES include/SDL include
+	PATH_SUFFIXES include/SDL2 include
 	PATHS ${SDL2TTF_SEARCH_PATHS}
 )
 

camera/gstCamera.cpp

@@ -29,22 +29,19 @@ uint32_t gstCamera::mSize   = 0;
 */
 bool gstCamera::mOnboardCamera = false;
 
-#define HEIGHT 720
-#define WIDTH 1280
 #define DEPTH 12
 #define SIZE HEIGHT * WIDET * DEPTH / 8
 
-
 // constructor
-gstCamera::gstCamera()
+gstCamera::gstCamera(int height, int width)
 {	
 	mAppSink    = NULL;
 	mBus        = NULL;
 	mPipeline   = NULL;	
 	mRGBA       = NULL;
 
-	mWidth     = 1280;
-	mHeight    = 720;
+	mWidth     = height;
+	mHeight    = width;
 	mDepth     = 12;
 	mSize      = (mWidth * mHeight * mDepth) / 8;
 
@@ -108,8 +105,8 @@ bool gstCamera::ConvertYUVtoRGBA( void* input, void** output )
 	}
 
 	// nvcamera is NV12
-//	if( CUDA_FAILED(cudaYUYVToRGBAf((uint8_t*)input, (float4*)mRGBA, mWidth, mHeight)) )
-//		return false;
+	if( CUDA_FAILED(cudaYUVToRGBAf((uint8_t*)input, (float4*)mRGBA, mWidth, mHeight)) )
+		return false;
 
 	*output = mRGBA;
 	return true;
@@ -332,25 +329,25 @@ gstCamera* gstCamera::Create()
 	std::ostringstream ss;
 
 	ss << "nvcamerasrc fpsRange=\"30.0 30.0\" ! video/x-raw(memory:NVMM), width=(int)" 
-		<< WIDTH 
+		<< 1280 
 		<< ", height=(int)" 
-		<< HEIGHT 
+		<< 720 
 		<< ", format=(string)NV12 ! nvvidconv flip-method=2 ! "; 
 	ss << "video/x-raw ! appsink name=mysink";
         mOnboardCamera = true;
-        return Create(ss.str());
+        return Create(ss.str(), 720, 1280);
 }
 
 // Create
-gstCamera* gstCamera::Create(std::string pipeline)
+gstCamera* gstCamera::Create(std::string pipeline, int height, int width)
 {
 	if( !gstreamerInit() )
 	{
 		printf(LOG_GSTREAMER "failed to initialize gstreamer API\n");
 		return NULL;
 	}
 
-	gstCamera* cam = new gstCamera();
+	gstCamera* cam = new gstCamera(width, height);
 
 	if( !cam )
 		return NULL;

camera/gstCamera.h

@@ -21,7 +21,9 @@ class gstCamera
 {
 public:
 	static gstCamera* Create();
-	static gstCamera* Create(std::string pipeline);
+	static gstCamera* Create(std::string pipeline, int height, int width);
+
+	gstCamera(int height, int width);
 	~gstCamera();
 
 	bool Open();

cuda/cudaRGB-NV12.cu

@@ -33,7 +33,7 @@ cudaError_t cudaRGBToRGBAf( uint8_t* srcDev, float4* destDev, size_t width, size
 	if( !srcDev || !destDev )
 		return cudaErrorInvalidDevicePointer;
 
-	const dim3 blockDim(8,8,1);
+	const dim3 blockDim(8,1,1);
 	const dim3 gridDim(width/blockDim.x, height/blockDim.y, 1);
 
 	RGBToRGBAf<<<gridDim, blockDim>>>( (uint8_t*)srcDev, destDev, width, height );

cuda/cudaYUV-NV12.cu

@@ -19,7 +19,7 @@
 __constant__ uint32_t constAlpha;
 __constant__ float  constHueColorSpaceMat[9];
 
-
+#define LIMIT_RGB(x)    (((x)<0)?0:((x)>255)?255:(x))
 
 __device__ void YUV2RGB(uint32_t *yuvi, float *red, float *green, float *blue)
 {
@@ -48,6 +48,8 @@ __device__ void YUV2RGB(uint32_t *yuvi, float *red, float *green, float *blue)
 	const float u    = float(yuvi[1]) - 512.0f;
 	const float v    = float(yuvi[2]) - 512.0f;
 
+//printf("YUV luma=%f, u=%f, v=%f\n", luma, u, v);
+
    /*R = Y + 1.140V
    G = Y - 0.395U - 0.581V
    B = Y + 2.032U*/
@@ -59,9 +61,33 @@ __device__ void YUV2RGB(uint32_t *yuvi, float *red, float *green, float *blue)
 	*red    = luma + 1.140f * v;
 	*green  = luma - 0.395f * u - 0.581f * v;
 	*blue   = luma + 2.032f * u;
+//printf("RGB luma=%f, u=%f, v=%f\n", *red , *green, *blue);
 }
 
 
+__device__ void YUV82RGB(uint8_t *yuvi, float *red, float *green, float *blue)
+{
+	const float y0 = float(yuvi[0]);
+	const float cb0    = float(yuvi[1]);
+	const float cr0    = float(yuvi[2]);
+
+//printf("YUV luma=%f, u=%f, v=%f\n", y0, cb0, cr0);
+
+	unsigned char r0(LIMIT_RGB(round((298.082 * y0) / 256.0 + (408.583 * cr0) / 256.0 - 222.921)));
+
+
+	unsigned char g0(LIMIT_RGB(round((298.082 * y0) / 256.0 - (100.291 * cb0) / 256.0
+			- (208.120 * cr0) / 256.0 + 135.576)));
+
+
+	unsigned char b0(LIMIT_RGB(round((298.082 * y0) / 256.0 + (516.412 * cb0) / 256.0 - 276.836)));
+
+	*red    = (float)r0;
+	*green  = (float)g0;
+	*blue   = (float)b0;
+//printf("RGB red=%f, green=%f, blue=%f\n", *red , *green, *blue);
+}
+
 __device__ uint32_t RGBAPACK_8bit(float red, float green, float blue, uint32_t alpha)
 {
     uint32_t ARGBpixel = 0;
@@ -372,6 +398,81 @@ cudaError_t cudaNV12ToRGBAf( uint8_t* srcDev, float4* destDev, size_t width, siz
 }
 
 
+//-------------------------------------------------------------------------------------------------------------------------
+// RTP YUV color space conversion
+
+__global__ void YUVToRGBAf(uint32_t* srcImage,  size_t nSourcePitch,
+                           float4* dstImage,     size_t nDestPitch,
+                           uint32_t width,       uint32_t height)
+{
+    int x, y;
+    uint32_t yuv101010Pel[2];
+    uint32_t processingPitch = ((width) + 63) & ~63;
+    uint8_t *srcImageU8     = (uint8_t *)srcImage;
+
+    processingPitch = nSourcePitch * 2;
+
+    // Pad borders with duplicate pixels, and we multiply by 2 because we process 2 pixels per thread
+    x = blockIdx.x * (blockDim.x << 1) + (threadIdx.x << 1);
+    y = blockIdx.y *  blockDim.y       +  threadIdx.y;
+
+    if (x >= width)
+        return; //x = width - 1;
+
+    if (y >= height)
+        return; // y = height - 1;
+
+//printf(">>>>> processingPitch %u, x %u, y %u \n", processingPitch, x, y );
+
+    // this steps performs the color conversion
+    uint8_t yuvi[6];
+    float red[2], green[2], blue[2];
+
+    yuvi[0] = srcImageU8[y * processingPitch + x*2 +1]; //Y0
+    yuvi[1] = srcImageU8[y * processingPitch + x*2 ];//Cb
+    yuvi[2] = srcImageU8[y * processingPitch + x*2 + 2];//Cr
+
+    yuvi[3] = srcImageU8[y * processingPitch + x*2 + 3];//Y1
+    yuvi[4] = srcImageU8[y * processingPitch + x*2];//Cb
+    yuvi[5] = srcImageU8[y * processingPitch + x*2 + 2];//Cr
+
+    // YUV to RGB Transformation conversion
+    YUV82RGB(&yuvi[0], &red[0], &green[0], &blue[0]);
+    YUV82RGB(&yuvi[3], &red[1], &green[1], &blue[1]);
+
+	dstImage[y * width + x]     = make_float4(red[0], green[0], blue[0], 1.0f);
+	dstImage[y * width + x + 1] = make_float4(red[1], green[1], blue[1], 1.0f);
+}
+
+
+
+// cudaNV12ToRGBA
+cudaError_t cudaYUVToRGBAf( uint8_t* srcDev, size_t srcPitch, float4* destDev, size_t destPitch, size_t width, size_t height )
+{
+	if( !srcDev || !destDev )
+		return cudaErrorInvalidDevicePointer;
+
+	if( srcPitch == 0 || destPitch == 0 || width == 0 || height == 0 )
+		return cudaErrorInvalidValue;
+
+	if( !nv12ColorspaceSetup )
+		cudaNV12SetupColorspace();
+
+	const dim3 blockDim(8,8,1);
+	//const dim3 gridDim((width+(2*blockDim.x-1))/(2*blockDim.x), (height+(blockDim.y-1))/blockDim.y, 1);
+	const dim3 gridDim(iDivUp(width,blockDim.x), iDivUp(height, blockDim.y), 1);
+
+	YUVToRGBAf<<<gridDim, blockDim>>>( (uint32_t*)srcDev, srcPitch, destDev, destPitch, width, height );
+
+	return CUDA(cudaGetLastError());
+}
+
+cudaError_t cudaYUVToRGBAf( uint8_t* srcDev, float4* destDev, size_t width, size_t height )
+{
+	return cudaYUVToRGBAf(srcDev, width * sizeof(uint8_t), destDev, width * sizeof(float4), width, height);
+}
+
+
 // cudaNV12SetupColorspace
 cudaError_t cudaNV12SetupColorspace( float hue )
 {

cuda/cudaYUV.h

@@ -122,6 +122,13 @@ cudaError_t cudaNV12ToRGBAf( uint8_t* input, float4* output, size_t width, size_
  */
 cudaError_t cudaNV12SetupColorspace( float hue = 0.0f ); 
 
+//////////////////////////////////////////////////////////////////////////////////
+/// @name YUV to RGBf
+//////////////////////////////////////////////////////////////////////////////////
+
+cudaError_t cudaYUVToRGBAf( uint8_t* input, size_t inputPitch, float4* output, size_t outputPitch, size_t width, size_t height );
+cudaError_t cudaYUVToRGBAf( uint8_t* input, float4* output, size_t width, size_t height );
+
 ///@}
 
 #endif

detectnet-camera/detectnet-camera.cpp

@@ -7,8 +7,14 @@
 #define GST_RTP_SRC 0
 #define SDL_DISPLAY 1
 #define ABACO 1
+
+#if GST_RTP_SRC
+#define HEIGHT 480
+#define WIDTH 640
+#else
 #define HEIGHT 720
 #define WIDTH 1280
+#endif
 
 #if V4L_CAMERA
 #include "v4l2Camera.h"
@@ -31,7 +37,6 @@
 
 #include "cudaMappedMemory.h"
 #include "cudaNormalize.h"
-#include "cudaFont.h"
 
 #include "detectNet.h"
 
@@ -91,7 +96,7 @@ int main( int argc, char** argv )
 	int height    = HEIGHT;
     std::ostringstream pipeline;
 #if GST_RTP_SRC
-	pipeline << "udpsrc port=5004 caps=\"application/x-rtp, media=(string)video, clock-rate=(int)90000, encoding-name=(string)RAW, sampling=(string)YCbCr-4:2:2, depth=(string)8, width=(string)" << width << ", height=(string)" << height << ", payload=(int)96\" ! ";
+	pipeline << "udpsrc address=239.192.1.44 port=5004 caps=\"application/x-rtp, media=(string)video, clock-rate=(int)90000, encoding-name=(string)RAW, sampling=(string)YCbCr-4:2:2, depth=(string)8, width=(string)" << width << ", height=(string)" << height << ", payload=(int)96\" ! ";
 	pipeline << "queue ! rtpvrawdepay ! queue ! ";
 	pipeline << "appsink name=mysink";
 #else
@@ -103,7 +108,7 @@ int main( int argc, char** argv )
 
     static  std::string pip = pipeline.str();
 
-	gstCamera* camera = gstCamera::Create(pip);
+	gstCamera* camera = gstCamera::Create(pip, height, width);
 #else
 	gstCamera* camera = gstCamera::Create();
 #endif
@@ -168,11 +173,6 @@ int main( int argc, char** argv )
 			printf("detectnet-camera:  failed to create openGL texture\n");
 	}
 
-	/*
-	 * create font
-	 */
-	cudaFont* font = cudaFont::Create();
-
 	/*
 	 * load logo
 	 */
@@ -248,15 +248,6 @@ int main( int argc, char** argv )
 					lastStart = n;
 				}
 			}
-
-			/*if( font != NULL )
-			{
-				char str[256];
-				sprintf(str, "%05.2f%% %s", confidence * 100.0f, net->GetClassDesc(img_class));
-				
-				font->RenderOverlay((float4*)imgRGBA, (float4*)imgRGBA, camera->GetWidth(), camera->GetHeight(),
-								    str, 10, 10, make_float4(255.0f, 255.0f, 255.0f, 255.0f));
-			}*/
 
 			if( display != NULL )
 			{
@@ -299,9 +290,12 @@ int main( int argc, char** argv )
 				texture->Render(0,0);		
 			}
 #if ABACO
-            texture->RenderText("WE INNOVATE. WE DELIVER. ", black, camera->GetWidth()-381, camera->GetHeight()-33, 18); 
-            texture->RenderText("YOU SUCCEED.", white, camera->GetWidth()-140, camera->GetHeight()-33, 18); 
-            texture->RenderText("abaco.com", white, 15, camera->GetHeight()-40, 28); 
+			if (!(camera->GetWidth() < 1024))
+			{
+				texture->RenderText("WE INNOVATE. WE DELIVER. ", black, camera->GetWidth()-381, camera->GetHeight()-33, 18); 
+				texture->RenderText("YOU SUCCEED.", white, camera->GetWidth()-140, camera->GetHeight()-33, 18); 
+				texture->RenderText("abaco.com", white, 15, camera->GetHeight()-40, 28); 
+			}
 #endif
 			display->EndRender();
 		}