Feature/mono8 cuda conversion

camera/gstCamera.cpp

@@ -62,6 +62,7 @@ gstCamera::gstCamera()
 	mDepth  = 0;
 	mSize   = 0;
 	mSource = GST_SOURCE_NVCAMERA;
+	mInverted = false;
 
 	mLatestRGBA       = 0;
 	mLatestRingbuffer = 0;
@@ -411,10 +412,10 @@ bool gstCamera::buildLaunchStr( gstCameraSrc src )
 		mSource = src;	 // store camera source method
 
 	#if NV_TENSORRT_MAJOR > 1 && NV_TENSORRT_MAJOR < 5	// if JetPack 3.1-3.3 (different flip-method)
-		const int flipMethod = 0;					// Xavier (w/TRT5) camera is mounted inverted
+		const int flipMethod = mInverted ? 2 : 0;					// Xavier (w/TRT5) camera is mounted inverted
 	#else
-		const int flipMethod = 2;
-	#endif	
+		const int flipMethod = mInverted ? 0 : 2;
+	#endif
 
 		if( src == GST_SOURCE_NVCAMERA )
 			ss << "nvcamerasrc fpsRange=\"30.0 30.0\" ! video/x-raw(memory:NVMM), width=(int)" << mWidth << ", height=(int)" << mHeight << ", format=(string)NV12 ! nvvidconv flip-method=" << flipMethod << " ! "; //'video/x-raw(memory:NVMM), width=(int)1920, height=(int)1080, format=(string)I420, framerate=(fraction)30/1' ! ";
@@ -483,7 +484,7 @@ bool gstCamera::parseCameraStr( const char* camera )
 
 
 // Create
-gstCamera* gstCamera::Create( uint32_t width, uint32_t height, const char* camera )
+gstCamera* gstCamera::Create( uint32_t width, uint32_t height, const char* camera, bool inverted)
 {
 	if( !gstreamerInit() )
 	{
@@ -503,6 +504,7 @@ gstCamera* gstCamera::Create( uint32_t width, uint32_t height, const char* camer
 	cam->mHeight     = height;
 	cam->mDepth      = cam->csiCamera() ? 12 : 24;	// NV12 or RGB
 	cam->mSize       = (width * height * cam->mDepth) / 8;
+	cam->mInverted   = inverted;
 
 	if( !cam->init(GST_SOURCE_NVARGUS) )
 	{

camera/gstCamera.h

@@ -111,7 +111,7 @@ class gstCamera
 	 *
 	 * @returns A pointer to the created gstCamera device, or NULL if there was an error.
 	 */
-	static gstCamera* Create( uint32_t width, uint32_t height, const char* camera=NULL );
+	static gstCamera* Create( uint32_t width, uint32_t height, const char* camera=NULL, bool inverted=false );
 
 	/**
 	 * Release the camera interface and resources.
@@ -292,6 +292,7 @@ class gstCamera
 	uint32_t mHeight;
 	uint32_t mDepth;
 	uint32_t mSize;
+	bool mInverted;
 
 	static const uint32_t NUM_RINGBUFFERS = 16;
 

cuda/cudaRGB.cu

@@ -239,5 +239,49 @@ cudaError_t cudaRGBA32ToBGR8( float4* srcDev, uchar3* destDev, size_t width, siz
 	return cudaRGBA32ToBGR8(srcDev, destDev, width, height, make_float2(0.0f, 255.0f));
 }
 
+//-------------------------------------------------------------------------------------------------------------------------
+__global__ void RGBAToMONO8(float4* srcImage,
+							unsigned char* dstImage,
+                           int width, int height,
+					  	   float scaling_factor)
+{
+	const int x = (blockIdx.x * blockDim.x) + threadIdx.x;
+	const int y = (blockIdx.y * blockDim.y) + threadIdx.y;
+
+	const int pixel = y * width + x;
+
+	if( x >= width )
+		return; 
 
+	if( y >= height )
+		return;
+
+	const float4 px = srcImage[pixel];
+	const float val = ( px.x + px.y + px.z) * scaling_factor / 3;
+
+	dstImage[pixel] = (unsigned char)val;
+}
+
+cudaError_t cudaRGBA32ToMONO8( float4* srcDev, unsigned char* destDev, size_t width, size_t height, const float2& inputRange )
+{
+	if( !srcDev || !destDev )
+		return cudaErrorInvalidDevicePointer;
+
+	if( width == 0 || height == 0 )
+		return cudaErrorInvalidValue;
+
+	const float multiplier = 255.0f / inputRange.y;
+
+	const dim3 blockDim(8,8,1);
+	const dim3 gridDim(iDivUp(width,blockDim.x), iDivUp(height,blockDim.y), 1);
+
+	RGBAToMONO8<<<gridDim, blockDim>>>( srcDev, destDev, width, height, multiplier );
+
+	return CUDA(cudaGetLastError());
+}
+
+cudaError_t cudaRGBA32ToMONO8( float4* srcDev, unsigned char* destDev, size_t width, size_t height )
+{
+	return cudaRGBA32ToMONO8(srcDev, destDev, width, height, make_float2(0.0f, 255.0f));
+}
 

cuda/cudaRGB.h

@@ -123,5 +123,24 @@ cudaError_t cudaRGBA32ToBGRA8( float4* input, uchar4* output, size_t width, size
 
 ///@}
 
+
+//////////////////////////////////////////////////////////////////////////////////
+/// @name Floating-point RGBA to 8-bit grayscale
+/// @ingroup colorspace
+//////////////////////////////////////////////////////////////////////////////////
+
+///@{
+
+/**
+ * Convert 32-bit floating-point RGBA image into 8-bit fixed-point MONO8 image.
+ * Assumes 0.0-255.0f input range, output range is 0-255.
+ * @ingroup colorspace
+ */
+cudaError_t cudaRGBA32ToMONO8( float4* srcDev, unsigned char* destDev, size_t width, size_t height, const float2& inputRange );
+
+cudaError_t cudaRGBA32ToMONO8( float4* srcDev, unsigned char* destDev, size_t width, size_t height );
+
+///@}
+
 #endif