updated CUDA colorspace conversions

dusty-nv · Jun 17, 2020 · 16056b7 · 16056b7
1 parent ec68ca5
commit 16056b7
Show file tree

Hide file tree

Showing 5 changed files with 76 additions and 21 deletions.
diff --git a/cuda/cudaColorspace.cpp b/cuda/cudaColorspace.cpp
@@ -302,17 +302,25 @@ cudaError_t cudaConvertColor( void* input, imageFormat inputFormat,
 			return CUDA(cudaGray8ToRGB32((uint8_t*)input, (float3*)output, width, height));
 		else if( outputFormat == IMAGE_RGBA32F || outputFormat == IMAGE_BGRA32F )
 			return CUDA(cudaGray8ToRGBA32((uint8_t*)input, (float4*)output, width, height));
+		else if( outputFormat == IMAGE_GRAY8 )
+			return CUDA(cudaMemcpy(output, input, imageFormatSize(inputFormat, width, height), cudaMemcpyDeviceToDevice));
+		else if( outputFormat == IMAGE_GRAY32F )
+			return CUDA(cudaGray8ToGray32((uint8_t*)input, (float*)output, width, height));
 	}
 	else if( inputFormat == IMAGE_GRAY32F )
 	{
 		if( outputFormat == IMAGE_RGB8 || outputFormat == IMAGE_BGR8 )
-			return CUDA(cudaGray32ToRGB8((float*)input, (uchar3*)output, width, height));
+			return CUDA(cudaGray32ToRGB8((float*)input, (uchar3*)output, width, height, pixel_range));
 		else if( outputFormat == IMAGE_RGBA8 || outputFormat == IMAGE_BGRA8 )
-			return CUDA(cudaGray32ToRGBA8((float*)input, (uchar4*)output, width, height));
+			return CUDA(cudaGray32ToRGBA8((float*)input, (uchar4*)output, width, height, pixel_range));
 		else if( outputFormat == IMAGE_RGB32F || outputFormat == IMAGE_BGR32F )
 			return CUDA(cudaGray32ToRGB32((float*)input, (float3*)output, width, height));
 		else if( outputFormat == IMAGE_RGBA32F || outputFormat == IMAGE_BGRA32F )
 			return CUDA(cudaGray32ToRGBA32((float*)input, (float4*)output, width, height));
+		else if( outputFormat == IMAGE_GRAY8 )
+			return CUDA(cudaGray32ToGray8((float*)input, (uint8_t*)output, width, height, pixel_range));
+		else if( outputFormat == IMAGE_GRAY32F )
+			return CUDA(cudaMemcpy(output, input, imageFormatSize(inputFormat, width, height), cudaMemcpyDeviceToDevice));
 	}
 	else if( imageFormatIsBayer(inputFormat) )
 	{

diff --git a/cuda/cudaGrayscale.cu b/cuda/cudaGrayscale.cu
@@ -266,6 +266,12 @@ cudaError_t cudaGray32ToRGBA32( float* srcDev, float4* dstDev, size_t width, siz
 	return launchGrayToRGB<float, float4>(srcDev, dstDev, width, height);
 }
 
+// cudaGray8ToGray32 (uint8 -> float)
+cudaError_t cudaGray8ToGray32( uint8_t* srcDev, float* dstDev, size_t width, size_t height )
+{
+	return launchGrayToRGB<uint8_t, float>(srcDev, dstDev, width, height);
+}
+
 
 //-----------------------------------------------------------------------------------
 // Grayscale to RGB (normalized)
@@ -318,4 +324,13 @@ cudaError_t cudaGray32ToRGB8( float* srcDev, uchar3* dstDev, size_t width, size_
 cudaError_t cudaGray32ToRGBA8( float* srcDev, uchar4* dstDev, size_t width, size_t height, const float2& inputRange )
 {
 	return launchGrayToRGB_Norm<float, uchar4>(srcDev, dstDev, width, height, inputRange);
-}
+}
+
+// cudaGray32ToGray8 (float -> uint8)
+cudaError_t cudaGray32ToGray8( float* srcDev, uint8_t* dstDev, size_t width, size_t height, const float2& inputRange )
+{
+	return launchGrayToRGB_Norm<float, uint8_t>(srcDev, dstDev, width, height, inputRange);
+}
+
+
+
diff --git a/cuda/cudaGrayscale.h b/cuda/cudaGrayscale.h
@@ -27,9 +27,39 @@
 #include "cudaUtility.h"
 
 
+//////////////////////////////////////////////////////////////////////////////////
+/// @name 8-bit grayscale to floating-point grayscale (and vice versa)
+/// @see cudaConvertColor() from cudaColorspace.h for automated format conversion
+/// @ingroup colorspace
+//////////////////////////////////////////////////////////////////////////////////
+
+///@{
+
+/**
+ * Convert uint8 grayscale image into float grayscale.
+ * @ingroup colorspace
+ */
+cudaError_t cudaGray8ToGray32( uint8_t* input, float* output, size_t width, size_t height );
+
+/**
+ * Convert float grayscale image into uint8 grayscale.
+ *
+ * @param pixelRange specifies the floating-point pixel value range of the input image, 
+ *                   which is used to rescale the fixed-point pixel outputs to [0,255].
+ *                   The default input range is [0,255], where no rescaling occurs.
+ *                   Other common input ranges are [-1, 1] or [0,1].
+ *
+ * @ingroup colorspace
+ */
+cudaError_t cudaGray32ToGray8( float* input, uint8_t* output, size_t width, size_t height, 
+						 const float2& pixelRange=make_float2(0,255) );
+
+///@}
+
+
 //////////////////////////////////////////////////////////////////////////////////
 /// @name RGB/BGR to 8-bit grayscale
-/// @see cudaConvertColor() from cudaColorspace.h for automated conversion
+/// @see cudaConvertColor() from cudaColorspace.h for automated format conversion
 /// @ingroup colorspace
 //////////////////////////////////////////////////////////////////////////////////
 
@@ -92,7 +122,7 @@ cudaError_t cudaRGBA32ToGray8( float4* input, uint8_t* output, size_t width, siz
 
 //////////////////////////////////////////////////////////////////////////////////
 /// @name RGB/BGR to floating-point grayscale
-/// @see cudaConvertColor() from cudaColorspace.h for automated conversion
+/// @see cudaConvertColor() from cudaColorspace.h for automated format conversion
 /// @ingroup colorspace
 //////////////////////////////////////////////////////////////////////////////////
 
@@ -143,7 +173,7 @@ cudaError_t cudaRGBA32ToGray32( float4* input, float* output, size_t width, size
 
 //////////////////////////////////////////////////////////////////////////////////
 /// @name 8-bit grayscale to RGB/BGR
-/// @see cudaConvertColor() from cudaColorspace.h for automated conversion
+/// @see cudaConvertColor() from cudaColorspace.h for automated format conversion
 /// @ingroup colorspace
 //////////////////////////////////////////////////////////////////////////////////
 
@@ -178,7 +208,7 @@ cudaError_t cudaGray8ToRGBA32( uint8_t* input, float4* output, size_t width, siz
 
 //////////////////////////////////////////////////////////////////////////////////
 /// @name Floating-point grayscale to RGB/BGR
-/// @see cudaConvertColor() from cudaColorspace.h for automated conversion
+/// @see cudaConvertColor() from cudaColorspace.h for automated format conversion
 /// @ingroup colorspace
 //////////////////////////////////////////////////////////////////////////////////
 
@@ -195,7 +225,8 @@ cudaError_t cudaGray8ToRGBA32( uint8_t* input, float4* output, size_t width, siz
  *
  * @ingroup colorspace
  */
-cudaError_t cudaGray32ToRGB8( float* input, uchar3* output, size_t width, size_t height, const float2& pixelRange=make_float2(0,255) );
+cudaError_t cudaGray32ToRGB8( float* input, uchar3* output, size_t width, size_t height, 
+						const float2& pixelRange=make_float2(0,255) );
 
 /**
  * Convert float grayscale image into uchar4 RGB/BGR.
@@ -207,7 +238,8 @@ cudaError_t cudaGray32ToRGB8( float* input, uchar3* output, size_t width, size_t
  *
  * @ingroup colorspace
  */
-cudaError_t cudaGray32ToRGBA8( float* input, uchar4* output, size_t width, size_t height, const float2& pixelRange=make_float2(0,255) );
+cudaError_t cudaGray32ToRGBA8( float* input, uchar4* output, size_t width, size_t height, 
+						 const float2& pixelRange=make_float2(0,255) );
 
 /**
  * Convert float grayscale image into float3 RGB/BGR.

diff --git a/cuda/cudaRGB.h b/cuda/cudaRGB.h
@@ -30,7 +30,7 @@
 
 //////////////////////////////////////////////////////////////////////////////////
 /// @name RGB/RGBA to BGR/BGRA (or vice-versa)
-/// @see cudaConvertColor() from cudaColorspace.h for automated conversion
+/// @see cudaConvertColor() from cudaColorspace.h for automated format conversion
 /// @ingroup colorspace
 //////////////////////////////////////////////////////////////////////////////////
 
@@ -73,7 +73,7 @@ cudaError_t cudaRGBA32ToBGRA32( float4* input, float4* output, size_t width, siz
 
 //////////////////////////////////////////////////////////////////////////////////
 /// @name 8-bit RGB/BGR to 8-bit RGBA/BGRA (or vice-versa)
-/// @see cudaConvertColor() from cudaColorspace.h for automated conversion
+/// @see cudaConvertColor() from cudaColorspace.h for automated format conversion
 /// @ingroup colorspace
 //////////////////////////////////////////////////////////////////////////////////
 
@@ -106,7 +106,7 @@ cudaError_t cudaRGBA8ToRGB8( uchar4* input, uchar3* output, size_t width, size_t
 
 //////////////////////////////////////////////////////////////////////////////////
 /// @name Floating-point RGB/BGR to floating-point RGBA/BGRA (or vice versa)
-/// @see cudaConvertColor() from cudaColorspace.h for automated conversion
+/// @see cudaConvertColor() from cudaColorspace.h for automated format conversion
 /// @ingroup colorspace
 //////////////////////////////////////////////////////////////////////////////////
 
@@ -139,7 +139,7 @@ cudaError_t cudaRGBA32ToRGB32( float4* input, float3* output, size_t width, size
 
 //////////////////////////////////////////////////////////////////////////////////
 /// @name 8-bit images to floating-point images
-/// @see cudaConvertColor() from cudaColorspace.h for automated conversion
+/// @see cudaConvertColor() from cudaColorspace.h for automated format conversion
 /// @ingroup colorspace
 //////////////////////////////////////////////////////////////////////////////////
 
@@ -194,7 +194,7 @@ cudaError_t cudaRGBA8ToRGBA32( uchar4* input, float4* output, size_t width, size
 
 //////////////////////////////////////////////////////////////////////////////////
 /// @name Floating-point images to 8-bit images
-/// @see cudaConvertColor() from cudaColorspace.h for automated conversion
+/// @see cudaConvertColor() from cudaColorspace.h for automated format conversion
 /// @ingroup colorspace
 //////////////////////////////////////////////////////////////////////////////////
 

diff --git a/cuda/cudaYUV.h b/cuda/cudaYUV.h
@@ -29,7 +29,7 @@
 
 //////////////////////////////////////////////////////////////////////////////////
 /// @name YUV I420 4:2:0 planar to RGB
-/// @see cudaConvertColor() from cudaColorspace.h for automated conversion
+/// @see cudaConvertColor() from cudaColorspace.h for automated format conversion
 /// @ingroup colorspace
 //////////////////////////////////////////////////////////////////////////////////
 
@@ -64,7 +64,7 @@ cudaError_t cudaI420ToRGBA(void* input, float4* output, size_t width, size_t hei
 
 //////////////////////////////////////////////////////////////////////////////////
 /// @name YUV YV12 4:2:0 planar to RGB
-/// @see cudaConvertColor() from cudaColorspace.h for automated conversion
+/// @see cudaConvertColor() from cudaColorspace.h for automated format conversion
 /// @ingroup colorspace
 //////////////////////////////////////////////////////////////////////////////////
 
@@ -99,7 +99,7 @@ cudaError_t cudaYV12ToRGBA(void* input, float4* output, size_t width, size_t hei
 
 //////////////////////////////////////////////////////////////////////////////////
 /// @name RGB to YUV I420 4:2:0 planar
-/// @see cudaConvertColor() from cudaColorspace.h for automated conversion
+/// @see cudaConvertColor() from cudaColorspace.h for automated format conversion
 /// @ingroup colorspace
 //////////////////////////////////////////////////////////////////////////////////
 
@@ -134,7 +134,7 @@ cudaError_t cudaRGBAToI420( float4* input, void* output, size_t width, size_t he
 
 //////////////////////////////////////////////////////////////////////////////////
 /// @name RGB to YUV YV12 4:2:0 planar
-/// @see cudaConvertColor() from cudaColorspace.h for automated conversion
+/// @see cudaConvertColor() from cudaColorspace.h for automated format conversion
 /// @ingroup colorspace
 //////////////////////////////////////////////////////////////////////////////////
 
@@ -169,7 +169,7 @@ cudaError_t cudaRGBAToYV12( float4* input, void* output, size_t width, size_t he
 
 //////////////////////////////////////////////////////////////////////////////////
 /// @name YUV YUYV 4:2:2 packed to RGB
-/// @see cudaConvertColor() from cudaColorspace.h for automated conversion
+/// @see cudaConvertColor() from cudaColorspace.h for automated format conversion
 /// @ingroup colorspace
 //////////////////////////////////////////////////////////////////////////////////
 
@@ -204,7 +204,7 @@ cudaError_t cudaYUYVToRGBA( void* input, float4* output, size_t width, size_t he
 
 //////////////////////////////////////////////////////////////////////////////////
 /// @name YUV UYVY 4:2:2 packed to RGB
-/// @see cudaConvertColor() from cudaColorspace.h for automated conversion
+/// @see cudaConvertColor() from cudaColorspace.h for automated format conversion
 /// @ingroup colorspace
 //////////////////////////////////////////////////////////////////////////////////
 
@@ -239,7 +239,7 @@ cudaError_t cudaUYVYToRGBA( void* input, float4* output, size_t width, size_t he
 
 //////////////////////////////////////////////////////////////////////////////////
 /// @name YUV NV12 to RGB
-/// @see cudaConvertColor() from cudaColorspace.h for automated conversion
+/// @see cudaConvertColor() from cudaColorspace.h for automated format conversion
 /// @ingroup colorspace
 //////////////////////////////////////////////////////////////////////////////////