CUDA 10.0 使用 cudaBindTexture2D 示例代码

下面是一个使用 cudaBindTexture2D 的示例代码：

#include <cuda_runtime.h>
#include <stdio.h>

texture<float, cudaTextureType2D> tex;

__global__ void kernel(float *out, int width, int height)
{
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < width && y < height) {
        out[y * width + x] = tex2D(tex, x, y);
    }
}

int main()
{
    int width = 512;
    int height = 512;
    float *data = new float[width * height];
    for (int i = 0; i < width * height; i++) {
        data[i] = i;
    }
    
    cudaArray *cuArray;
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
    cudaMallocArray(&cuArray, &channelDesc, width, height);
    cudaMemcpyToArray(cuArray, 0, 0, data, width * height * sizeof(float), cudaMemcpyHostToDevice);

    cudaBindTextureToArray(tex, cuArray, channelDesc);

    int block_size = 16;
    dim3 dimBlock(block_size, block_size);
    dim3 dimGrid((width + block_size - 1) / block_size, (height + block_size - 1) / block_size);

    float *out;
    cudaMalloc(&out, width * height * sizeof(float));

    kernel<<<dimGrid, dimBlock>>>(out, width, height);

    cudaUnbindTexture(tex);

    cudaFreeArray(cuArray);
    cudaFree(out);
    delete[] data;

    return 0;
}

在这个示例中，我们首先创建了一个大小为 512 x 512 的 float 数组，然后将其拷贝到了一个 cudaArray 中。接着，我们使用 cudaBindTextureToArray 将这个 cudaArray 绑定到了一个 cuda 纹理 tex 上。最后，我们使用 kernel 函数将 tex 中的数据读取到了一个大小为 512 x 512 的 float 数组中，并将其输出。最后，我们使用 cudaUnbindTexture 解绑了纹理，并释放了所有分配的内存。