CUDA の Texture Memory で float3 のような3要素の型は fetch できない

試した環境

本題

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);

参考

docs.nvidia.com