This repository has been archived by the owner on Nov 8, 2021. It is now read-only.
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
Convert RGBA texture to NV12 CUDA buffer for Windows 7.
- Loading branch information
1 parent
c3173f8
commit 7540851
Showing
8 changed files
with
230 additions
and
20 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,90 @@ | ||
| <?xml version="1.0" encoding="utf-8"?> | ||
| <Project DefaultTargets="Build" ToolsVersion="15.0" xmlns="http://schemas.microsoft.com/developer/msbuild/2003"> | ||
| <ItemGroup Label="ProjectConfigurations"> | ||
| <ProjectConfiguration Include="Debug|x64"> | ||
| <Configuration>Debug</Configuration> | ||
| <Platform>x64</Platform> | ||
| </ProjectConfiguration> | ||
| <ProjectConfiguration Include="Release|x64"> | ||
| <Configuration>Release</Configuration> | ||
| <Platform>x64</Platform> | ||
| </ProjectConfiguration> | ||
| </ItemGroup> | ||
| <PropertyGroup Label="Globals"> | ||
| <ProjectGuid>{12336478-A663-42F8-8D52-8D093CD99992}</ProjectGuid> | ||
| <RootNamespace>CUDA</RootNamespace> | ||
| <WindowsTargetPlatformVersion>10.0.17134.0</WindowsTargetPlatformVersion> | ||
| </PropertyGroup> | ||
| <Import Project="$(VCTargetsPath)\Microsoft.Cpp.Default.props" /> | ||
| <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug|x64'" Label="Configuration"> | ||
| <ConfigurationType>StaticLibrary</ConfigurationType> | ||
| <UseDebugLibraries>true</UseDebugLibraries> | ||
| <CharacterSet>MultiByte</CharacterSet> | ||
| <PlatformToolset>v141</PlatformToolset> | ||
| </PropertyGroup> | ||
| <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|x64'" Label="Configuration"> | ||
| <ConfigurationType>StaticLibrary</ConfigurationType> | ||
| <UseDebugLibraries>false</UseDebugLibraries> | ||
| <WholeProgramOptimization>true</WholeProgramOptimization> | ||
| <CharacterSet>MultiByte</CharacterSet> | ||
| <PlatformToolset>v141</PlatformToolset> | ||
| </PropertyGroup> | ||
| <Import Project="$(VCTargetsPath)\Microsoft.Cpp.props" /> | ||
| <ImportGroup Label="ExtensionSettings"> | ||
| <Import Project="$(VCTargetsPath)\BuildCustomizations\CUDA 9.2.props" /> | ||
| </ImportGroup> | ||
| <ImportGroup Label="PropertySheets" Condition="'$(Configuration)|$(Platform)'=='Debug|x64'"> | ||
| <Import Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" /> | ||
| </ImportGroup> | ||
| <ImportGroup Label="PropertySheets" Condition="'$(Configuration)|$(Platform)'=='Release|x64'"> | ||
| <Import Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" /> | ||
| </ImportGroup> | ||
| <PropertyGroup Label="UserMacros" /> | ||
| <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug|x64'"> | ||
| <LinkIncremental>true</LinkIncremental> | ||
| </PropertyGroup> | ||
| <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Debug|x64'"> | ||
| <ClCompile> | ||
| <WarningLevel>Level3</WarningLevel> | ||
| <Optimization>Disabled</Optimization> | ||
| <PreprocessorDefinitions>WIN32;WIN64;_DEBUG;_CONSOLE;%(PreprocessorDefinitions)</PreprocessorDefinitions> | ||
| </ClCompile> | ||
| <Link> | ||
| <GenerateDebugInformation>true</GenerateDebugInformation> | ||
| <SubSystem>Console</SubSystem> | ||
| <AdditionalDependencies>cudart_static.lib;kernel32.lib;user32.lib;gdi32.lib;winspool.lib;comdlg32.lib;advapi32.lib;shell32.lib;ole32.lib;oleaut32.lib;uuid.lib;odbc32.lib;odbccp32.lib;%(AdditionalDependencies)</AdditionalDependencies> | ||
| </Link> | ||
| <CudaCompile> | ||
| <TargetMachinePlatform>64</TargetMachinePlatform> | ||
| </CudaCompile> | ||
| </ItemDefinitionGroup> | ||
| <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release|x64'"> | ||
| <ClCompile> | ||
| <WarningLevel>Level3</WarningLevel> | ||
| <Optimization>MaxSpeed</Optimization> | ||
| <FunctionLevelLinking>true</FunctionLevelLinking> | ||
| <IntrinsicFunctions>true</IntrinsicFunctions> | ||
| <PreprocessorDefinitions>WIN32;WIN64;NDEBUG;_CONSOLE;%(PreprocessorDefinitions)</PreprocessorDefinitions> | ||
| </ClCompile> | ||
| <Link> | ||
| <GenerateDebugInformation>true</GenerateDebugInformation> | ||
| <EnableCOMDATFolding>true</EnableCOMDATFolding> | ||
| <OptimizeReferences>true</OptimizeReferences> | ||
| <SubSystem>Console</SubSystem> | ||
| <AdditionalDependencies>cudart_static.lib;kernel32.lib;user32.lib;gdi32.lib;winspool.lib;comdlg32.lib;advapi32.lib;shell32.lib;ole32.lib;oleaut32.lib;uuid.lib;odbc32.lib;odbccp32.lib;%(AdditionalDependencies)</AdditionalDependencies> | ||
| </Link> | ||
| <CudaCompile> | ||
| <TargetMachinePlatform>64</TargetMachinePlatform> | ||
| </CudaCompile> | ||
| </ItemDefinitionGroup> | ||
| <ItemGroup> | ||
| <CudaCompile Include="RGBToNV12.cu" /> | ||
| </ItemGroup> | ||
| <ItemGroup> | ||
| <ClInclude Include="RGBToNV12.h" /> | ||
| </ItemGroup> | ||
| <Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" /> | ||
| <ImportGroup Label="ExtensionTargets"> | ||
| <Import Project="$(VCTargetsPath)\BuildCustomizations\CUDA 9.2.targets" /> | ||
| </ImportGroup> | ||
| </Project> |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,88 @@ | ||
| #include <stdint.h> | ||
|
|
||
| #include "RGBToNV12.h" | ||
|
|
||
| __device__ float rgb2y(uchar4 c) { | ||
| return 0.257f * c.x + 0.504f * c.y + 0.098f * c.z + 16.0f; | ||
| } | ||
|
|
||
| __device__ float rgb2u(uchar4 c) { | ||
| return -0.148f * c.x - 0.291f * c.y + 0.439f * c.z + 128.0f; | ||
| } | ||
|
|
||
| __device__ float rgb2v(uchar4 c) { | ||
| return 0.439f * c.x - 0.368f * c.y - 0.071f * c.z + 128.0f; | ||
| } | ||
|
|
||
| texture<uchar4, cudaTextureType2D, cudaReadModeElementType> texRef; | ||
|
|
||
| __global__ void RGBA2NV12_kernel(uint8_t *dstImage, size_t destPitch, | ||
| uint32_t width, uint32_t height) | ||
| { | ||
| // Pad borders with duplicate pixels, and we multiply by 2 because we process 2 pixels per thread | ||
| int32_t x = blockIdx.x * (blockDim.x << 1) + (threadIdx.x << 1); | ||
| int32_t y = blockIdx.y * (blockDim.y << 1) + (threadIdx.y << 1); | ||
|
|
||
| int x1 = x + 1; | ||
| int y1 = y + 1; | ||
|
|
||
| if (x1 >= width) | ||
| return; //x = width - 1; | ||
|
|
||
| if (y1 >= height) | ||
| return; // y = height - 1; | ||
|
|
||
| uchar4 c00 = tex2D(texRef, x, y); | ||
| uchar4 c01 = tex2D(texRef, x1, y); | ||
| uchar4 c10 = tex2D(texRef, x, y1); | ||
| uchar4 c11 = tex2D(texRef, x1, y1); | ||
|
|
||
| uint8_t y00 = (uint8_t)(rgb2y(c00) + 0.5f); | ||
| uint8_t y01 = (uint8_t)(rgb2y(c01) + 0.5f); | ||
| uint8_t y10 = (uint8_t)(rgb2y(c10) + 0.5f); | ||
| uint8_t y11 = (uint8_t)(rgb2y(c11) + 0.5f); | ||
|
|
||
| uint8_t u = (uint8_t)((rgb2u(c00) + rgb2u(c01) + rgb2u(c10) + rgb2u(c11)) * 0.25f + 0.5f); | ||
| uint8_t v = (uint8_t)((rgb2v(c00) + rgb2v(c01) + rgb2v(c10) + rgb2v(c11)) * 0.25f + 0.5f); | ||
|
|
||
| dstImage[destPitch * y + x] = y00; | ||
| dstImage[destPitch * y + x1] = y01; | ||
| dstImage[destPitch * y1 + x] = y10; | ||
| dstImage[destPitch * y1 + x1] = y11; | ||
|
|
||
| uint32_t chromaOffset = destPitch * height; | ||
| int32_t x_chroma = x; | ||
| int32_t y_chroma = y >> 1; | ||
|
|
||
| dstImage[chromaOffset + destPitch * y_chroma + x_chroma] = u; | ||
| dstImage[chromaOffset + destPitch * y_chroma + x_chroma + 1] = v; | ||
| } | ||
|
|
||
| extern "C" | ||
| cudaError_t RGBA2NV12(cudaArray *srcImage, | ||
| uint8_t *dstImage, size_t destPitch, | ||
| uint32_t width, uint32_t height) | ||
| { | ||
| cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsigned); | ||
|
|
||
| // Set texture parameters | ||
| texRef.addressMode[0] = cudaAddressModeWrap; | ||
| texRef.addressMode[1] = cudaAddressModeWrap; | ||
| texRef.filterMode = cudaFilterModePoint; | ||
| texRef.normalized = false; | ||
|
|
||
| cudaError_t cudaStatus = cudaBindTextureToArray(texRef, srcImage, channelDesc); | ||
| if (cudaStatus != cudaSuccess) { | ||
| return cudaStatus; | ||
| } | ||
|
|
||
| dim3 block(32, 16, 1); | ||
| dim3 grid((width + (2 * block.x - 1)) / (2 * block.x), (height + (2 * block.y - 1)) / (2 * block.y), 1); | ||
|
|
||
| RGBA2NV12_kernel<<<grid, block>>>(dstImage, destPitch, width, height); | ||
|
|
||
| cudaThreadSynchronize(); | ||
|
|
||
| cudaStatus = cudaGetLastError(); | ||
| return cudaStatus; | ||
| } |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,8 @@ | ||
| #pragma once | ||
|
|
||
| #include <cuda.h> | ||
|
|
||
| extern "C" | ||
| cudaError_t RGBA2NV12(cudaArray *srcImage, | ||
| uint8_t *dstImage, size_t destPitch, | ||
| uint32_t width, uint32_t height); |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters