cuda - 如何将 CUDA tex1DFetch 与 cudaTextureObject_t 一起使用?
问题描述
当我注意到纹理引用已被弃用时,我正在使用它们,我尝试更新我的测试函数以使用带有 tex1Dfetch 的“新”无绑定纹理对象,但无法产生相同的结果。
我目前正在探索使用纹理内存来加速我的 aho-corasick 实现;我能够tex1D()
使用纹理引用,但是,我注意到它们已被弃用并决定使用纹理对象。
当我尝试以任何方式使用结果时,我的内核出现了一些非常奇怪的行为;我可以做到results[tidx] = tidx;
没有任何问题,但results[tidx] = temp + 1;
只返回temp
nottemp * 3
或任何其他涉及temp
.
我看不出这种行为的逻辑原因,并且文档示例看起来很相似,以至于我看不出哪里出错了。
我已经阅读了 CUDA tex1Dfetch() 错误行为和 New CUDA Texture Object — 在 2D 情况下获取错误数据,但似乎都与我遇到的问题无关。
以防万一它有所作为;我正在使用带有 Nvidia GTX 980ti 的 CUDA 版本 10.0、V10.0.130。
#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
有谁知道到底出了什么问题?
解决方案
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 = {};
推荐阅读
- php - Symfony2,如何编辑表单中的字段并避免空密码字段
- python - 网络摄像头不通过 OpenCV 读取,但通过 guvcview 读取
- c# - 如何在包含 jpg 文件的每个目录和子目录中创建文件夹?
- javascript - 在firebase中更新值后数组不会重置
- azure - 如何替换当前变量的内容?
- javascript - 是否可以更改 ap 中特定值的样式?
- json - Groovy:有效地处理嵌套的 JSON 数组对象
- caching - Joomla 3.8 网站无法清除缓存
- php - Laravel - Eloquent:API 资源 - 条件关系
- apache-spark - ValueError:feature_columns 的项目必须是 _FeatureColumn。(张量流 1.13)