試した環境
- Microsoft Visual Studio Community 2019 Version 16.8.6
- NVIDIA CUDA 11.0
本題
CUDA の Texture Memory は int4 や float2 のような 1,2,4 要素の整数型と単精度浮動小数点型が利用できます。 逆に言うと float3 のような3要素の型は使えません。
https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#texture-memory
The type of a texel, which is restricted to the basic integer and single-precision floating-point types and any of the 1-, 2-, and 4-component vector types defined in Built-in Vector Types that are derived from the basic integer and single-precision floating-point types.
以下、 Texture Object を利用した例です。 基本型である float も含めて、以下の4つのカーネル関数はビルドが通ります。
__global__ void addKernel(float *dst, cudaTextureObject_t tex_src) { int i = threadIdx.x; float src = tex1Dfetch<float>(tex_src, i); dst[i] = src + 1.0f; }
__global__ void addKernel(float *dst, cudaTextureObject_t tex_src) { int i = threadIdx.x; float1 src = tex1Dfetch<float1>(tex_src, i); dst[i] = src.x + 1.0f; }
__global__ void addKernel(float *dst, cudaTextureObject_t tex_src) { int i = threadIdx.x; float2 src = tex1Dfetch<float2>(tex_src, i); dst[i] = src.x + src.y; }
__global__ void addKernel(float *dst, cudaTextureObject_t tex_src) { int i = threadIdx.x; float4 src = tex1Dfetch<float4>(tex_src, i); dst[i] = src.x + src.y + src.z + src.w; }
一方以下のfloat3 を使った例ではビルドエラーになります。
__global__ void addKernel(float *dst, cudaTextureObject_t tex_src) { int i = threadIdx.x; float3 src = tex1Dfetch<float3>(tex_src, i); dst[i] = src.x + src.y + src.z; }
no instance of overloaded function "tex1Dfetch" matches the argument list CudaTextureObject
float2 の場合の Texture Object の設定は以下の通りです。
cudaCreateChannelDesc(32, 32, 0, 0, cudaChannelFormatKindFloat)
とします。
cudaCreateChannelDesc(64, 0, 0, 0, cudaChannelFormatKindFloat)
などとするとcudaCreateTextureObject
で失敗します。
cudaResourceDesc resDesc; resDesc.resType = cudaResourceTypeLinear; resDesc.res.linear.devPtr = dev_src; resDesc.res.linear.desc = cudaCreateChannelDesc(32, 32, 0, 0, cudaChannelFormatKindFloat); resDesc.res.linear.sizeInBytes = size * sizeof(float2); cudaTextureDesc texDesc; std::memset(&texDesc, 0, sizeof(texDesc)); texDesc.readMode = cudaReadModeElementType; cudaTextureObject_t tex_src = 0; cudaCreateTextureObject(&tex_src, &resDesc, &texDesc, NULL);