CUDA tex1Dfetch() 错误的行为
CUDA tex1Dfetch() wrong behaviour
我对 CUDA 编程很陌生,我面临着一个让我发疯的问题。这是怎么回事:我有非常简单的程序(仅用于学习目的),其中创建了一个输入图像和一个输出图像 16x16。输入图像初始化为 0..255 之间的值,然后绑定到纹理。CUDA 内核只是将输入图像复制到输出图像。输入图像值是通过调用 tex1Dfetch() 获得的,在某些情况下,它返回非常奇怪的值。请参阅下面的代码、内核内部的注释和程序的输出。该代码是完整且可编译的,因此您可以在 VC 中创建 CUDA 项目并将代码粘贴到主".cu"文件中。
请帮帮我!我做错了什么?
我正在使用VS 2013社区和CUDA SDK 6.5 + CUDA集成用于VS 2013。
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
texture<unsigned char> tex;
cudaError_t testMyKernel(unsigned char * inputImg, unsigned char * outputImg, int width, int height);
__global__ void myKernel(unsigned char *outImg, int width)
{
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
int idx = row*width + col;
__shared__ unsigned char input;
__shared__ unsigned char input2;
unsigned char *outPix = outImg + idx;
//It fetches strange value, for example, when the idx==0 then the input is 51.
//But I expect that input==idx (according to the input image initialization).
input = tex1Dfetch(tex, idx);
printf("Fetched for idx=%d: %dn", idx, input);
*outPix = input;
//Very strange is that when I test the following code then the tex1Dfetch() returns correct values.
if (idx == 0)
{
printf("nKernel test print:n");
for (int i = 0; i < 256; i++)
{
input2 = tex1Dfetch(tex, i);
printf("%d,", input2);
}
}
}
int main()
{
const int width = 16;
const int height = 16;
const int count = width * height;
unsigned char imgIn[count];
unsigned char imgOut[count];
for (int i = 0; i < count; i++)
{
imgIn[i] = i;
}
cudaError_t cudaStatus = testMyKernel(imgIn, imgOut, width, height);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "testMyKernel failed!");
return 1;
}
printf("nnOutput values:n");
for (int i = 0; i < height; i++)
{
for (int j = 0; j < width; j++)
{
printf("%d,", imgOut[i * width + j]);
}
}
printf("n");
cudaStatus = cudaDeviceReset();
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaDeviceReset failed!");
return 1;
}
getchar();
return 0;
}
cudaError_t testMyKernel(unsigned char * inputImg, unsigned char * outputImg, int width, int height)
{
unsigned char * dev_in;
unsigned char * dev_out;
size_t size = width * height * sizeof(unsigned char);
cudaError_t cudaStatus;
cudaStatus = cudaSetDevice(0);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?");
goto Error;
}
// input data
cudaStatus = cudaMalloc((void**)&dev_in, size);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
cudaStatus = cudaMemcpy(dev_in, inputImg, size, cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
cudaStatus = cudaBindTexture(NULL, tex, dev_in, size);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaBindTexture failed!");
goto Error;
}
// output data
cudaStatus = cudaMalloc((void**)&dev_out, size);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
dim3 threadsPerBlock(4, 4);
int blk_x = width / threadsPerBlock.x;
int blk_y = height / threadsPerBlock.y;
dim3 numBlocks(blk_x, blk_y);
// Launch a kernel on the GPU with one thread for each element.
myKernel<<<numBlocks, threadsPerBlock>>>(dev_out, width);
cudaStatus = cudaGetLastError();
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "myKernel launch failed: %sn", cudaGetErrorString(cudaStatus));
goto Error;
}
cudaStatus = cudaDeviceSynchronize();
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching myKernel!n", cudaStatus);
goto Error;
}
//copy output image to host
cudaStatus = cudaMemcpy(outputImg, dev_out, size, cudaMemcpyDeviceToHost);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
Error:
cudaUnbindTexture(tex);
cudaFree(dev_in);
cudaFree(dev_out);
return cudaStatus;
}
这是程序的输出(截断一点):
Fetched for idx=0: 51
Fetched for idx=1: 51
Fetched for idx=2: 51
Fetched for idx=3: 51
Fetched for idx=16: 51
Fetched for idx=17: 51
Fetched for idx=18: 51
Fetched for idx=19: 51
Fetched for idx=32: 51
Fetched for idx=33: 51
Fetched for idx=34: 51
Fetched for idx=35: 51
Fetched for idx=48: 51
Fetched for idx=49: 51
Fetched for idx=50: 51
Fetched for idx=51: 51
Fetched for idx=192: 243
Fetched for idx=193: 243
Fetched for idx=194: 243
Fetched for idx=195: 243
Fetched for idx=208: 243
Fetched for idx=209: 243
Fetched for idx=210: 243
Fetched for idx=211: 243
Fetched for idx=224: 243
etc... (output truncated.. see the Output values)
Kernel test print:
0,1,2,3,4,5,6,7,8,9,10,11,12,13,14,15,16,17,18,19,20,21,22,23,24,25,26,27,28,29,
30,31,32,33,34,35,36,37,38,39,40,41,42,43,44,45,46,47,48,49,50,51,52,53,54,55,56
etc...(correct values)
Output values:
51,51,51,51,55,55,55,55,59,59,59,59,63,63,63,63,51,51,51,51,55,55,55,55,59,59,59
,59,63,63,63,63,51,51,51,51,55,55,55,55,59,59,59,59,63,63,63,63,51,51,51,51,55,55,
etc.. (wrong values)
内核
的这一行
input = tex1Dfetch(tex, idx);
导致块的线程之间出现争用情况。块中的所有线程都试图将值从纹理获取到__shared__
变量中,input
同时导致未定义的行为。您应该以__shared__
数组的形式为块的每个线程分配单独的共享内存空间。
对于您目前的情况,它可能是这样的
__shared__ unsigned char input[16]; //4 x 4 block size
内核的其余部分应如下所示:
int idx_local = threadIdx.y * blockDim.x + threadIdx.x; //local id of thread in a block
input[idx_local] = tex1Dfetch(tex, idx);
printf("Fetched for idx=%d: %dn", idx, input[idx_local]);
*outPix = input[idx_local];
内核末尾条件中的代码工作正常,因为由于指定的条件if (idx == 0)
,只有第一个块的第一个线程将串行执行所有处理,而所有其他线程将保持空闲状态,因此问题将由于没有竞争条件而消失。
相关文章:
- 警告处理为错误这里有什么问题
- "error: no matching function for call to"构造函数错误
- boost::进程间消息队列引发错误
- C++,OpenCV,尝试显示图像时"OpenCV(4.3.0) Error: Assertion failed (size.width>0 && size.height>0)"此错误
- 有关插入适配器的错误。[错误]请求从 'back_insert_iterator<vector<>>' 类型转换为非标量类型
- QT在错误的班级中寻找空位
- vector.resize()中的分配错误
- 代码在main()中运行,但在函数中出现错误
- 释放错误后堆使用
- (C++)分析树以计算返回错误值的简单算术表达式
- Project Euler问题4的错误解决方案
- 我的字符计数代码计算错误.为什么
- 从"int*"强制转换为"unsigned int"会丢失精度错误
- 尝试导入pybind-opencv模块时出现libgtk错误
- CMake项目Boost库错误:Boost/config/compiler/gcc.hpp:165:10:致命错误:cs
- 在某些循环内使用vector.push_back时出现分段错误
- MSVC多行宏编译器错误
- 静态数据成员的问题-修复链接错误会导致编译器错误
- 为什么在运行时没有向我们提供有关分段错误的更多信息?
- CUDA tex1Dfetch() 错误的行为