CUDAでテクスチャ

前回の記事では CUDA の線形メモリ（Linear Memory, cudaMalloc() で生成）を使って FP16 浮動小数点数値を格納しました。
今回は CUDA の Array と Texture を使って、0x0000 ~ 0xFFFF の範囲の 16bit 整数値（ushort）を正規化して 0.0 ~ 1.0 の範囲の 32bit 浮動小数点数値（float）として取り出すということをやってみます。ただし今回は半精度浮動小数点数（half）型は使いません。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <cstdio>
#include <cstdint>
#include <conio.h>

cudaError_t addWithCuda(float *c, const uint16_t *a, const uint16_t *b, size_t size);

texture<uint16_t, cudaTextureType1D, cudaReadModeNormalizedFloat> texA;
texture<uint16_t, cudaTextureType1D, cudaReadModeNormalizedFloat> texB;

__global__ void addKernel(float *c)
{
    const int i = threadIdx.x;
    const float texCoord = (float)i;
    const float a = tex1D(texA, texCoord);
    const float b = tex1D(texB, texCoord);
    c[i] = a + b;
}

int main()
{
    const int arraySize = 8;
    // 16bit 整数値を正規化して 32bit 浮動小数点数に変換するので、
    // 最初から 32bit の int や float で入力する場合と比べてどうしても精度は劣る。
    const uint16_t a[arraySize] = { 0xFFFF, 0x0000, 0x8000, 0x8000, 0xFFFF, 0x2000, 0x1000, 0x0000 };
    const uint16_t b[arraySize] = { 0x0000, 0xFFFF, 0x4000, 0x0000, 0xFFFF, 0x8000, 0xEFFF, 0x0000 };
    float c[arraySize] = { 0 };

    // Add vectors in parallel.
    cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "addWithCuda failed!");
        return 1;
    }

    for (int i = 0; i < arraySize; ++i)
    {
        printf("%f + %f = %f\n", a[i] / double(0xFFFF), b[i] / double(0xFFFF), c[i]);
    }

    // cudaDeviceReset must be called before exiting in order for profiling and
    // tracing tools such as Nsight and Visual Profiler to show complete traces.
    cudaStatus = cudaDeviceReset();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceReset failed!");
        return 1;
    }

    puts("Press any...");
    _getch();

    return 0;
}

// Helper function for using CUDA to add vectors in parallel.
cudaError_t addWithCuda(float *c, const uint16_t *a, const uint16_t *b, size_t size)
{
    cudaArray_t arrayA = nullptr;
    cudaArray_t arrayB = nullptr;
    decltype(c) dev_c = nullptr;

    const auto descForHalf = cudaCreateChannelDesc<uint16_t>();

    cudaError_t cudaStatus;

    // Choose which GPU to run on, change this on a multi-GPU system.
    cudaStatus = cudaSetDevice(0);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
        goto Error;
    }

    // Allocate GPU buffers for three vectors (two input, one output).
    cudaStatus = cudaMallocArray(&arrayA, &descForHalf, size);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMallocArray failed!");
        goto Error;
    }

    cudaStatus = cudaMallocArray(&arrayB, &descForHalf, size);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMallocArray failed!");
        goto Error;
    }

    cudaStatus = cudaMalloc(&dev_c, size * sizeof(decltype(*c)));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    // Copy input vectors from host memory to GPU buffers.
    cudaStatus = cudaMemcpyToArray(arrayA, 0, 0, a, size * sizeof(uint16_t), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpyToArray failed!");
        goto Error;
    }

    cudaStatus = cudaMemcpyToArray(arrayB, 0, 0, b, size * sizeof(uint16_t), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpyToArray failed!");
        goto Error;
    }

    cudaBindTextureToArray(texA, arrayA, descForHalf);
    cudaBindTextureToArray(texB, arrayB, descForHalf);

    // Launch a kernel on the GPU with one thread for each element.
    addKernel<<<1, static_cast<uint32_t>(size)>>>(dev_c);

    // cudaDeviceSynchronize waits for the kernel to finish, and returns
    // any errors encountered during the launch.
    cudaStatus = cudaDeviceSynchronize();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
        goto Error;
    }

    // Copy output vector from GPU buffer to host memory.
    cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(decltype(*c)), cudaMemcpyDeviceToHost);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

Error:
    cudaFreeArray(arrayA);
    cudaFreeArray(arrayB);
    cudaFree(dev_c);

    return cudaStatus;
}

コード中で重要なのは cudaReadModeNormalizedFloat です。
cudaReadModeNormalizedFloat を使うと、texture をサンプリングするときに [0, 1] もしくは [-1, 1] の範囲に正規化した浮動小数点数として取り出すことができます。
例えば
ushort であれば [0, 65535] が [0.0, 1.0] にリニアマッピングされます。
byte であれば [0, 255] が [0.0, 1.0] にリニアマッピングされます。
short であれば [-32768, 32767] が [-1.0, 1.0] にリニアマッピングされます。
sbyte であれば [-128, 127] が [-1.0, 1.0] にリニアマッピングされます。

OpenGL や Direct3D のシェーダーにおいて、チャンネルごとに 0-255 の固定レンジ 8bit 値を持つ R8G8B8A8 テクスチャをサンプリングするときに float4 として取得できるのによく似ています。
おそらくテクスチャフェッチするときにハードウェア機能で正規化するはずなので、カーネルコードでユーザーが明示的にスケーリングして正規化するよりも高速化できると思われます。

なお、CUDA には tex1D() の他にも tex1Dfetch() という読み出し関数があるのですが、tex1D() のほうは CUDA Array にバインドした Texture 用、tex1Dfetch() のほうは線形メモリにバインドした Texture 用らしいです。CUDA における Texture というのはリソースそのものではなく、リソースに対するビュー（DirectX 10/11 の SRV）と考えたほうがよさそうです。
ちなみに DirectX 11.0 フル対応の Fermi アーキテクチャ以降では、UAV に相当する Surface が使えます。