CUDA 10.0 使用 cudaBindTexture2D 绑定纹理并设置偏移量示例

以下是一个使用 cudaBindTexture2D 函数并带有 offset 参数的示例代码：

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

texture<float, cudaTextureType2D> tex;

__global__ void kernel(float* output, int width, int height, size_t pitch, int offset)
{
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    if (x < width && y < height)
    {
        float value = tex2D<float>(tex, x + offset, y);
        output[y * pitch / sizeof(float) + x] = value;
    }
}

int main()
{
    const int width = 640;
    const int height = 480;
    const int offset = 10;

    float* data = new float[width * height];
    for (int i = 0; i < width * height; i++)
    {
        data[i] = i;
    }

    cudaArray* array = nullptr;
    cudaMallocArray(&array, &tex.channelDesc, width + offset, height);
    cudaMemcpyToArray(array, 0, 0, data, width * height * sizeof(float), cudaMemcpyHostToDevice);

    cudaBindTexture2D(nullptr, &tex, array, &tex.channelDesc, width + offset, height, pitch);
    cudaError_t error = cudaGetLastError();
    if (error != cudaSuccess)
    {
        printf("cudaBindTexture2D failed with error: %s\n", cudaGetErrorString(error));
        return -1;
    }

    float* output = nullptr;
    cudaMalloc(&output, width * height * sizeof(float));
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks((width + threadsPerBlock.x - 1) / threadsPerBlock.x, (height + threadsPerBlock.y - 1) / threadsPerBlock.y);
    kernel<<<numBlocks, threadsPerBlock>>>(output, width, height, pitch, offset);

    cudaMemcpy(data, output, width * height * sizeof(float), cudaMemcpyDeviceToHost);
    cudaUnbindTexture(&tex);
    cudaFreeArray(array);
    cudaFree(output);

    delete[] data;
    return 0;
}

在此示例代码中，我们首先创建了一个大小为 width * height 的浮点数数组 data，并将其初始化为连续的 0 到 (width*height-1) 的值。

然后我们使用 cudaMallocArray 函数分配了一个大小为 (width+offset) * height 的 cudaArray 数组，并使用 cudaMemcpyToArray 函数将数据从主机端复制到设备端。

接下来，我们使用 cudaBindTexture2D 函数将 tex 纹理绑定到 cudaArray 数组上，并使用 offset 参数指定偏移量。

在核函数中，我们首先使用 tex2D 函数从纹理中读取一个浮点数值，并将其写入到输出数组中。请注意，我们使用了 offset 参数来访问纹理中的正确像素。

最后，我们使用 cudaMemcpy 函数将输出数组从设备端复制到主机端，并释放了所有的设备端内存。

请注意，此示例代码中的所有 CUDA 函数调用都没有进行错误检查。在实际应用中，我们应该始终检查 CUDA 函数的返回值，并在发生错误时采取适当的措施。