CUDA tex1Dfetch() 错误行为
CUDA tex1Dfetch() wrong behaviour
我是 CUDA 编程的新手，我正面临一个让我抓狂的问题。这是怎么回事：
我有一个非常简单的程序（仅用于学习目的），其中创建了一张输入图像和一张输出图像 16x16。输入图像被初始化为 0..255 之间的值，然后绑定到纹理。 CUDA 内核只是将输入图像复制到输出图像。输入图像值是通过调用 tex1Dfetch() 获得的，在某些情况下 returns 非常奇怪的值。请看下面的代码，内核中的注释和程序的输出。代码完整且可编译，因此您可以在 VC 中创建一个 CUDA 项目并将代码粘贴到主“.cu”文件中。
请帮帮我！我做错了什么？
我正在使用 VS 2013 社区和 CUDA SDK 6.5 + VS 2013 的 CUDA 集成。
#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: %d\n", 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("\n\nOutput 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: %s\n", 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__ 数组的形式为块的每个线程分配单独的共享内存 space。
对于您当前的情况，可能类似于
__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: %d\n", idx, input[idx_local]);
*outPix = input[idx_local];
内核末尾的条件内的代码工作正常，因为由于指定的条件 if (idx == 0)，只有第一个块的第一个线程会串行执行所有处理，而所有其他线程会保持空闲状态，因此问题将因不存在竞争条件而消失。
我是 CUDA 编程的新手，我正面临一个让我抓狂的问题。这是怎么回事： 我有一个非常简单的程序（仅用于学习目的），其中创建了一张输入图像和一张输出图像 16x16。输入图像被初始化为 0..255 之间的值，然后绑定到纹理。 CUDA 内核只是将输入图像复制到输出图像。输入图像值是通过调用 tex1Dfetch() 获得的，在某些情况下 returns 非常奇怪的值。请看下面的代码，内核中的注释和程序的输出。代码完整且可编译，因此您可以在 VC 中创建一个 CUDA 项目并将代码粘贴到主“.cu”文件中。
请帮帮我！我做错了什么？
我正在使用 VS 2013 社区和 CUDA SDK 6.5 + VS 2013 的 CUDA 集成。
#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: %d\n", 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("\n\nOutput 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: %s\n", 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__ 数组的形式为块的每个线程分配单独的共享内存 space。
对于您当前的情况，可能类似于
__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: %d\n", idx, input[idx_local]);
*outPix = input[idx_local];
内核末尾的条件内的代码工作正常，因为由于指定的条件 if (idx == 0)，只有第一个块的第一个线程会串行执行所有处理，而所有其他线程会保持空闲状态，因此问题将因不存在竞争条件而消失。