如何将 CUDA tex1DFetch 与 cudaTextureObject_t 一起使用？

Question

当我注意到它们已被弃用时，我正在使用纹理引用，我尝试更新我的测试函数以使用带有 tex1Dfetch 的 'new' 无绑定纹理对象，但无法产生相同的结果。

我目前正在探索使用纹理内存来加速我的 aho-corasick 实现；我能够 tex1D() 使用纹理引用，但是，我注意到它们已被弃用，因此决定改用纹理对象。

当我尝试以任何方式使用结果时，内核出现了一些非常奇怪的行为；我可以 results[tidx] = tidx; 没有任何问题，但是 results[tidx] = temp + 1; 只有 returns temp 的值而不是 temp * 3 或涉及 temp 的任何其他数值测试.

我看不出这种行为的逻辑原因，而且文档示例看起来非常相似，以至于我看不出哪里出错了。

我已经阅读了 CUDA tex1Dfetch() 错误行为和新 CUDA 纹理对象 — 在 2D 情况下获取错误数据，但似乎都与我遇到的问题无关。

以防万一我正在使用 CUDA 版本 10.0，V10.0.130 和 Nvidia GTX 980ti。

#include <iostream>

__global__ void test(cudaTextureObject_t tex ,int* results){
    int tidx = threadIdx.y * blockDim.x + threadIdx.x;
    unsigned temp = tex1Dfetch<unsigned>(tex, threadIdx.x);
    results[tidx] = temp * 3;
}

int main(){
    int *host_arr;
    const int host_arr_size = 8;

    // Create and populate host array
    std::cout << "Host:" << std::endl;
    cudaMallocHost(&host_arr, host_arr_size*sizeof(int));
    for (int i = 0; i < host_arr_size; ++i){
        host_arr[i] = i * 2;
        std::cout << host_arr[i] << std::endl;
    }

    // Create resource description
    struct cudaResourceDesc resDesc;
    resDesc.resType = cudaResourceTypeLinear;
    resDesc.res.linear.devPtr = &host_arr;
    resDesc.res.linear.sizeInBytes = host_arr_size*sizeof(unsigned);
    resDesc.res.linear.desc = cudaCreateChannelDesc<unsigned>();
    // Create texture description
    struct cudaTextureDesc texDesc;
    texDesc.readMode = cudaReadModeElementType;
    // Create texture
    cudaTextureObject_t tex;
    cudaCreateTextureObject(&tex, &resDesc, &texDesc, NULL);

    // Allocate results array
    int * result_arr;
    cudaMalloc(&result_arr, host_arr_size*sizeof(unsigned));

    // launch test kernel
    test<<<1, host_arr_size>>>(tex, result_arr);

    // fetch results
    std::cout << "Device:" << std::endl;
    cudaMemcpy(host_arr, result_arr, host_arr_size*sizeof(unsigned), cudaMemcpyDeviceToHost);
    // print results
    for (int i = 0; i < host_arr_size; ++i){
        std::cout << host_arr[i] << std::endl;
    }

    // Tidy Up
    cudaDestroyTextureObject(tex);
    cudaFreeHost(host_arr);
    cudaFree(result_arr);
}

我希望上面的工作类似于下面的工作（确实有效）：


texture<int, 1, cudaReadModeElementType> tex_ref;
cudaArray* cuda_array;

__global__ void test(int* results){
    const int tidx = threadIdx.x;
    results[tidx] = tex1D(tex_ref, tidx) * 3;
}

int main(){
    int *host_arr;
    int host_arr_size = 8;

    // Create and populate host array
    cudaMallocHost((void**)&host_arr, host_arr_size * sizeof(int));
    for (int i = 0; i < host_arr_size; ++i){
        host_arr[i] = i * 2;
        std::cout << host_arr[i] << std::endl;
    }

    // bind to texture
    cudaChannelFormatDesc cuDesc = cudaCreateChannelDesc <int >();
    cudaMallocArray(&cuda_array, &cuDesc, host_arr_size);
    cudaMemcpyToArray(cuda_array, 0, 0, host_arr , host_arr_size * sizeof(int), cudaMemcpyHostToDevice);
    cudaBindTextureToArray(tex_ref , cuda_array);
    // Allocate results array
    int * result_arr;
    cudaMalloc((void**)&result_arr, host_arr_size*sizeof(int));

    // launch kernel
    test<<<1, host_arr_size>>>(result_arr);

    // fetch results
    cudaMemcpy(host_arr, result_arr, host_arr_size * sizeof(int), cudaMemcpyDeviceToHost);
    // print results
    for (int i = 0; i < host_arr_size; ++i){
        std::cout << host_arr[i] << std::endl;
    }

    // Tidy Up
    cudaUnbindTexture(tex_ref);
    cudaFreeHost(host_arr);
    cudaFreeArray(cuda_array);
    cudaFree(result_arr);
}

预期结果：

Host:
0
2
4
6
8
10
12
14
Device:
0
6
12
18
24
30
36
42

实际结果：

Host:
0
2
4
6
8
10
12
14
Device:
0
2
4
6
8
10
12
14

有谁知道到底出了什么问题？

Answer 1

CUDA API 函数调用 return 错误代码。您要检查这些错误代码。特别是当某些事情明显出错时某处…

您使用同一个数组来存储初始数组数据以及从设备接收结果。您的内核启动失败并出现非法地址错误，因为您没有有效的纹理对象。您没有有效的纹理对象，因为创建纹理对象失败。内核启动后的第一个 API 调用是 cudaMemcpy() 以获取结果。由于在内核启动期间出现错误，cudaMemcpy() 将失败，returning 最近的错误而不是执行复制。结果，您的 host_arr 缓冲区的内容没有改变，您只是再次显示原始输入数据。

创建纹理对象失败的原因在 documentation（强调我的）中进行了解释：

If cudaResourceDesc::resType is set to cudaResourceTypeLinear, cudaResourceDesc::res::linear::devPtr must be set to a valid device pointer, that is aligned to cudaDeviceProp::textureAlignment. […]

纹理对象无法引用主机内存。您的代码中的问题在于：

resDesc.res.linear.devPtr = &host_arr;

您需要在设备内存中分配一个缓冲区，例如，使用 cudaMalloc()，将您的数据复制到那里，并创建一个引用该设备缓冲区的纹理对象。

此外，您的 texDesc 未正确初始化。在您的情况下，只需对其进行零初始化就足够了：

struct cudaTextureDesc texDesc = {};

如何将 CUDA tex1DFetch 与 cudaTextureObject_t 一起使用？

How to use CUDA tex1DFetch with cudaTextureObject_t?

textures

cuda