大家好,我刚接触cuda编程。我想写一个类似于cook的光线追踪程序,不过在读取纹理时遇到一些问题。运行时cuda提示cudaErrorInvalidTexture,我对比了几个例程依然没找出错误来。具体代码如下。
texture<float4, 2, cudaReadModeElementType> tex;
//read texture
void readTexture(char *filename, uint width, uint height) {
//Load ppm picture
unsigned char *h_data;
sdkLoadPPM4ub(filename, &h_data, &width, &height);
float4 *h_output = new float4[width*height];
for(int i = 0; i < width*height; i++) {
h_output[i].x = h_data[4*i];
h_output[i].y = h_data[4*i+1];
h_output[i].z = h_data[4*i+2];
h_output[i].w = h_data[4*i+3];
//cout << h_output[i].x << " " << h_output[i].y << " " << h_output[i].z << " " << h_output[i].w << endl;
}
// Copy to device memory some data located at address h_data in host memory
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float4>();
cudaArray *cuArray;
checkCudaErrors(cudaMallocArray(&cuArray, &channelDesc, width, height));
cudaMemcpyToArray(cuArray, 0, 0, h_output, sizeof(h_output),cudaMemcpyHostToDevice);
// Set texture parameters
tex.addressMode[0] = cudaAddressModeWrap;
tex.addressMode[1] = cudaAddressModeWrap;
tex.filterMode = cudaFilterModeLinear;
tex.normalized = true; // access with normalized texture coordinates
// Bind the array to the texture
checkCudaErrors(cudaBindTextureToArray(tex, cuArray, channelDesc));
}
PS:如果要调用15个纹理,是全局定义15个纹理比较好,还是在每个类内定义纹理对象比较好呢?
楼主您好,通过阅读您的代码,未能发现任何问题。
您的texture<>声明、cudaMallocArray(), 以及cudaBindTextureToArray(), 均正确无误。
您的最后的cudaBindTextureToArray()应该返回cudaSuccess的。而不是cudaErrorInvalidTexture。
因为本人水平有限,错误在所难免,只能给出上述观点了。
我发现我工程配置的问题应该是。。感谢版主!
不过我读入一张512*1024的颜色为黄色的图片,贴出来的是红色的图。。输出h_output发现都是205,205,205.是这段代码的问题吗。。。?
如果您是普通的4B/像素的图片,那么您不应该使用float4作为元素类型的。
(那将会是16B!)
使用uchar4之类的也许不错。
不好意思。。还得打扰。。
我发现纹理总是不显示,然后__device__里输出结果发现颜色都是0.。。。
printf(“tex: %.2f %.2f %.2f”, tex2D(tex, 0.1, 0.1).x, tex2D(tex, 0.1, 0.1).y, tex2D(tex, 0.1, 0.1).z);
请问我可能是哪里出了错误= =|||?
不知道。
以及你没有回答我上个帖子中的询问。
(也许人家只是uchar4, 外加cudaReadModeNormalizedFloat返回的float4, 而不您这样的直接定义元素问float4的)
错就是错,对就是对。你不认识我,你不必在乎面子。
既然是4B, 请火速改成uchar4。
以及我看你需要float4, 不要忘记最后cudaReadModeNormalizedFloat。
感谢配合工作。
我使用的是cook(光线追踪的桌球)的pool_1.ppm那张图,颜色的确是4B的。
我改了一下代码。
cudaArray *cuArray;
texture<uchar4, 2, cudaReadModeNormalizedFloat> tex;
//read texture
void readTexture(char *filename,const char *exe_path, uint width, uint height) {
char *imgPath = sdkFindFilePath(filename, exe_path);
//Load ppm picture
unsigned char *h_data = new unsigned char[width*height*4];
if (sdkLoadPPM4(imgPath, &h_data, &width, &height) == 0)
cout << "Error load pic." << endl;
uchar4 *h_output = new uchar4[width*height];
for(int i = 0; i < width*height; i++) {
h_output[i].x = h_data[4*i];
h_output[i].y = h_data[4*i+1];
h_output[i].z = h_data[4*i+2];
h_output[i].w = h_data[4*i+3];
}
//cout << h_output[0].x << " " << h_output[0].y << " " << h_output[0].z << " " << h_output[0].w << endl;
// Copy to device memory some data located at address h_data in host memory
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<uchar4>();
checkCudaErrors(cudaMallocArray(&cuArray, &channelDesc, width, height));
cudaMemcpyToArray(cuArray, 0, 0, h_output, sizeof(h_output),cudaMemcpyHostToDevice);
// Set texture parameters
tex.addressMode[0] = cudaAddressModeWrap;
tex.addressMode[1] = cudaAddressModeWrap;
tex.filterMode = cudaFilterModeLinear;
tex.normalized = true; // access with normalized texture coordinates
// Bind the array to the texture
checkCudaErrors(cudaBindTextureToArray(tex, cuArray, channelDesc));
}
system
10
cudaMemcpyToArray(cuArray, 0, 0, h_output, sizeof(h_output),cudaMemcpyHostToDevice);
这句sizeof(h_output),改成sizeof(uchar4)widthheight,不然你只copy了4B(或者8B,在64位系统)的数据