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

Question

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

\n\n

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

\n\n

当我尝试以任何方式使用结果时，我会遇到一些非常奇怪的内核行为；我可以results[tidx] = tidx;没有任何问题，但只返回notresults[tidx] = temp + 1;的值或涉及 的任何其他数值测试。temptemp * 3temp

\n\n

我看不出这种行为的逻辑原因，并且文档示例看起来非常相似，我看不出哪里出了问题。

\n\n

我已经读过 CUDA tex1Dfetch() 错误行为和 New CUDA 纹理对象 \xe2\x80\x94 在 2D 情况下获取错误数据，但这似乎都与我遇到的问题无关。

\n\n

以防万一它有所作为；我正在使用 CUDA 版本 10.0、V10.0.130 和 Nvidia GTX 980ti。

\n\n

#include <iostream>\n\n__global__ void test(cudaTextureObject_t tex ,int* results){\n    int tidx = threadIdx.y * blockDim.x + threadIdx.x;\n    unsigned temp = tex1Dfetch<unsigned>(tex, threadIdx.x);\n    results[tidx] = temp * 3;\n}\n\nint main(){\n    int *host_arr;\n    const int host_arr_size = 8;\n\n    // Create and populate host array\n    std::cout << "Host:" << std::endl;\n    cudaMallocHost(&host_arr, host_arr_size*sizeof(int));\n    for (int i = 0; i < host_arr_size; ++i){\n        host_arr[i] = i * 2;\n        std::cout << host_arr[i] << std::endl;\n    }\n\n    // Create resource description\n    struct cudaResourceDesc resDesc;\n    resDesc.resType = cudaResourceTypeLinear;\n    resDesc.res.linear.devPtr = &host_arr;\n    resDesc.res.linear.sizeInBytes = host_arr_size*sizeof(unsigned);\n    resDesc.res.linear.desc = cudaCreateChannelDesc<unsigned>();\n    // Create texture description\n    struct cudaTextureDesc texDesc;\n    texDesc.readMode = cudaReadModeElementType;\n    // Create texture\n    cudaTextureObject_t tex;\n    cudaCreateTextureObject(&tex, &resDesc, &texDesc, NULL);\n\n    // Allocate results array\n    int * result_arr;\n    cudaMalloc(&result_arr, host_arr_size*sizeof(unsigned));\n\n    // launch test kernel\n    test<<<1, host_arr_size>>>(tex, result_arr);\n\n    // fetch results\n    std::cout << "Device:" << std::endl;\n    cudaMemcpy(host_arr, result_arr, host_arr_size*sizeof(unsigned), cudaMemcpyDeviceToHost);\n    // print results\n    for (int i = 0; i < host_arr_size; ++i){\n        std::cout << host_arr[i] << std::endl;\n    }\n\n    // Tidy Up\n    cudaDestroyTextureObject(tex);\n    cudaFreeHost(host_arr);\n    cudaFree(result_arr);\n}\n

\n\n

我预计上面的工作原理与下面的工作类似（确实有效）：

\n\n

\ntexture<int, 1, cudaReadModeElementType> tex_ref;\ncudaArray* cuda_array;\n\n__global__ void test(int* results){\n    const int tidx = threadIdx.x;\n    results[tidx] = tex1D(tex_ref, tidx) * 3;\n}\n\nint main(){\n    int *host_arr;\n    int host_arr_size = 8;\n\n    // Create and populate host array\n    cudaMallocHost((void**)&host_arr, host_arr_size * sizeof(int));\n    for (int i = 0; i < host_arr_size; ++i){\n        host_arr[i] = i * 2;\n        std::cout << host_arr[i] << std::endl;\n    }\n\n    // bind to texture\n    cudaChannelFormatDesc cuDesc = cudaCreateChannelDesc <int >();\n    cudaMallocArray(&cuda_array, &cuDesc, host_arr_size);\n    cudaMemcpyToArray(cuda_array, 0, 0, host_arr , host_arr_size * sizeof(int), cudaMemcpyHostToDevice);\n    cudaBindTextureToArray(tex_ref , cuda_array);\n    // Allocate results array\n    int * result_arr;\n    cudaMalloc((void**)&result_arr, host_arr_size*sizeof(int));\n\n    // launch kernel\n    test<<<1, host_arr_size>>>(result_arr);\n\n    // fetch results\n    cudaMemcpy(host_arr, result_arr, host_arr_size * sizeof(int), cudaMemcpyDeviceToHost);\n    // print results\n    for (int i = 0; i < host_arr_size; ++i){\n        std::cout << host_arr[i] << std::endl;\n    }\n\n    // Tidy Up\n    cudaUnbindTexture(tex_ref);\n    cudaFreeHost(host_arr);\n    cudaFreeArray(cuda_array);\n    cudaFree(result_arr);\n}\n

\n\n

预期成绩：

\n\n

Host:\n0\n2\n4\n6\n8\n10\n12\n14\nDevice:\n0\n6\n12\n18\n24\n30\n36\n42\n

\n\n

实际结果：

\n\n

Host:\n0\n2\n4\n6\n8\n10\n12\n14\nDevice:\n0\n2\n4\n6\n8\n10\n12\n14\n

\n\n

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

\n

Answer 1

CUDA API 函数调用返回错误代码。您想要检查这些错误代码。尤其是当某处明显出现问题时……

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

文档中解释了创建纹理对象失败的原因（强调我的）：

如果 cudaResourceDesc::resType 设置为 cudaResourceTypeLinear，则 cudaResourceDesc::res::linear::devPtr 必须设置为有效的设备指针，该指针与 cudaDeviceProp::textureAlignment 对齐。[…]

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

resDesc.res.linear.devPtr = &host_arr;

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

此外，您的texDesc初始化不正确。在您的情况下，只需将其初始化为零就足够了：

struct cudaTextureDesc texDesc = {};