大约两年前,我写了一个内核,同时处理几个数字网格.出现了一些非常奇怪的行为,导致错误的结果.当利用printf() - 内核中的语句来查找错误时,bug就消失了.
由于截止日期限制,我保持这种方式,虽然最近我认为这是不合适的编码风格.所以我重新访问了我的内核并将其归结为您在下面看到的内容.
__launch_bounds__(672, 2)
__global__ void heisenkernel(float *d_u, float *d_r, float *d_du, int radius,
int numNodesPerGrid, int numBlocksPerSM, int numGridsPerSM, int numGrids)
{
__syncthreads();
int id_sm = blockIdx.x / numBlocksPerSM; // (arbitrary) ID of Streaming Multiprocessor (SM) this thread works upon - (constant over lifetime of thread)
int id_blockOnSM = blockIdx.x % numBlocksPerSM; // Block number on this specific SM - (constant over lifetime of thread)
int id_r = id_blockOnSM * (blockDim.x - 2*radius) + threadIdx.x - radius; // Grid point number this thread is to work upon - (constant over lifetime of thread)
int id_grid = id_sm * numGridsPerSM; // Grid ID this thread is to work upon - (not constant over lifetime of thread)
while(id_grid < numGridsPerSM * (id_sm + 1)) // this loops over numGridsPerSM grids
{
__syncthreads();
int id_numInArray = id_grid * numNodesPerGrid + id_r; // Entry in array this thread is responsible for (read and possibly write) - (not constant over lifetime of thread)
float uchange = 0.0f;
//uchange = 1.0f; // if this line is uncommented, results will be computed correctly ("Solution 1")
float du = 0.0f;
if((threadIdx.x > radius-1) && (threadIdx.x < blockDim.x - radius) && (id_r < numNodesPerGrid) && (id_grid < numGrids))
{
if (id_r == 0) // FO-forward difference
du = (d_u[id_numInArray+1] - d_u[id_numInArray])/(d_r[id_numInArray+1] - d_r[id_numInArray]);
else if (id_r == numNodesPerGrid - 1) // FO-rearward difference
du = (d_u[id_numInArray] - d_u[id_numInArray-1])/(d_r[id_numInArray] - d_r[id_numInArray-1]);
else if (id_r == 1 || id_r == numNodesPerGrid - 2) //SO-central difference
du = (d_u[id_numInArray+1] - d_u[id_numInArray-1])/(d_r[id_numInArray+1] - d_r[id_numInArray-1]);
else if(id_r > 1 && id_r < numNodesPerGrid - 2)
du = d_fourpoint_constant * ((d_u[id_numInArray+1] - d_u[id_numInArray-1])/(d_r[id_numInArray+1] - d_r[id_numInArray-1])) + (1-d_fourpoint_constant) * ((d_u[id_numInArray+2] - d_u[id_numInArray-2])/(d_r[id_numInArray+2] - d_r[id_numInArray-2]));
else
du = 0;
}
__syncthreads();
if((threadIdx.x > radius-1 && threadIdx.x < blockDim.x - radius) && (id_r < numNodesPerGrid) && (id_grid < numGrids))
{
d_u[ id_numInArray] = d_u[id_numInArray] * uchange; // if this line is commented out, results will be computed correctly ("Solution 2")
d_du[ id_numInArray] = du;
}
__syncthreads();
++id_grid;
}
该内核计算了许多数值1D网格在所有网格点处的某个值的导数.
需要考虑的事项:(见底部的完整代码库)
网格由1300个网格点组成
每个网格必须由两个块处理(由于内存/寄存器限制)
每个块依次在37个网格上工作(或者更好:网格一半,while循环处理它)
每个线程负责每个网格中的相同网格点
对于要计算的导数,线程需要访问来自四个下一个网格点的数据
为了使块彼此独立,引入网格上的小重叠(每个网格的网格点666,667,668,669由来自不同块的两个线程读取,尽管只有一个线程正在写入它们,这是问题发生的重叠)
由于沸腾过程,块两侧的两个线程不进行计算,原来它们负责将相应的网格值写入共享内存
栅格的值被存储在u_arr
,du_arr
和r_arr
(和它们相应的器件阵列d_u
,d_du
和d_r
).每个网格在每个阵列中占用1300个连续值.内核中的while循环为每个块迭代超过37个网格.
为了评估内核的工作原理,每个网格都使用完全相同的值进行初始化,因此确定性程序将为每个网格生成相同的结果.我的代码不会发生这种情况.
Heisenbug的怪异:
我将网格0的计算值与每个其他网格进行了比较,并且在重叠处存在差异(网格点666-669),尽管不一致.有些网格有正确的值,有些则没有.连续两次运行会将不同的网格标记为错误.首先想到的是,在这个重叠的两个线程尝试同时写入内存,虽然情况似乎并非如此(我检查....并重新检查).
注释或取消注释行或printf()
用于调试目的也将改变程序的结果:当"询问"负责网格点的线程时,他们告诉我一切都是正确的,并且它们实际上是正确的.一旦我强制一个线程打印出它的变量,它们就会被正确计算(更重要的是:存储).使用Nsight Eclipse进行调试也是如此.
Memcheck/Racecheck:
cuda-memcheck(memcheck和racecheck)报告没有内存/竞争条件问题,但即使使用其中一个工具也能够影响结果的正确性.Valgrind给出了一些警告,但我认为它们与CUDA API有关,我无法影响它,这似乎与我的问题无关.
(更新)
正如所指出的,cuda-memcheck --tool racecheck
仅适用于共享内存竞争条件,而手头的问题具有竞争条件d_u
,即全局内存.
测试环境:
原始内核已经在不同的CUDA设备上进行了测试,具有不同的计算能力(2.0,3.0和3.5),每个配置中都会出现错误(以某种形式或其他形式).
我的(主要)测试系统如下:
2 x GTX 460,在运行X-server的GPU和另一台运行的GPU上进行了测试
驱动程序版本:340.46
Cuda Toolkit 6.5
Linux Kernel 3.11.0-12-generic(Linux Mint 16 - Xfce)
解决方案的状态:
到目前为止,我很确定一些内存访问是罪魁祸首,可能是编译器的一些优化或使用未初始化的值,而且我显然不了解一些基本的CUDA范例.事实上,printf()
内核中的语句(通过一些黑魔法也必须利用设备和主机内存)和memcheck算法(cuda-memcheck和valgrind)会影响同一方向的优点.
我很抱歉这个有点复杂的内核,但是我尽可能地将原始内核和调用放到了最后,这就是我所知道的.到目前为止,我已经学会了欣赏这个问题,我期待着了解这里发生了什么.
在代码中标记了两个"解决方案",它们强制内核按预期工作.
(更新)如下面的正确答案中所述,我的代码的问题是线程块边界处的竞争条件.由于每个网格上有两个块,并且无法保证哪个块首先工作,因此导致下面列出的行为.它还解释了使用代码中提到的"解决方案1"时的正确结果,因为输入/输出值d_u
不会改变uchange = 1.0
.
简单的解决方案是将此内核拆分为两个内核,一个计算d_u
,另一个计算衍生d_du
.更希望只有一个内核调用而不是两个,尽管我不知道如何实现这一点-arch=sm_20
.有-arch=sm_35
一个人可能会使用动态并行来实现这一点,尽管第二次内核调用的开销可以忽略不计.
heisenbug.cu:
#include
#include
#include
const float r_sol = 6.955E8f;
__constant__ float d_fourpoint_constant = 0.2f;
__launch_bounds__(672, 2)
__global__ void heisenkernel(float *d_u, float *d_r, float *d_du, int radius,
int numNodesPerGrid, int numBlocksPerSM, int numGridsPerSM, int numGrids)
{
__syncthreads();
int id_sm = blockIdx.x / numBlocksPerSM; // (arbitrary) ID of Streaming Multiprocessor (SM) this thread works upon - (constant over lifetime of thread)
int id_blockOnSM = blockIdx.x % numBlocksPerSM; // Block number on this specific SM - (constant over lifetime of thread)
int id_r = id_blockOnSM * (blockDim.x - 2*radius) + threadIdx.x - radius; // Grid point number this thread is to work upon - (constant over lifetime of thread)
int id_grid = id_sm * numGridsPerSM; // Grid ID this thread is to work upon - (not constant over lifetime of thread)
while(id_grid < numGridsPerSM * (id_sm + 1)) // this loops over numGridsPerSM grids
{
__syncthreads();
int id_numInArray = id_grid * numNodesPerGrid + id_r; // Entry in array this thread is responsible for (read and possibly write) - (not constant over lifetime of thread)
float uchange = 0.0f;
//uchange = 1.0f; // if this line is uncommented, results will be computed correctly ("Solution 1")
float du = 0.0f;
if((threadIdx.x > radius-1) && (threadIdx.x < blockDim.x - radius) && (id_r < numNodesPerGrid) && (id_grid < numGrids))
{
if (id_r == 0) // FO-forward difference
du = (d_u[id_numInArray+1] - d_u[id_numInArray])/(d_r[id_numInArray+1] - d_r[id_numInArray]);
else if (id_r == numNodesPerGrid - 1) // FO-rearward difference
du = (d_u[id_numInArray] - d_u[id_numInArray-1])/(d_r[id_numInArray] - d_r[id_numInArray-1]);
else if (id_r == 1 || id_r == numNodesPerGrid - 2) //SO-central difference
du = (d_u[id_numInArray+1] - d_u[id_numInArray-1])/(d_r[id_numInArray+1] - d_r[id_numInArray-1]);
else if(id_r > 1 && id_r < numNodesPerGrid - 2)
du = d_fourpoint_constant * ((d_u[id_numInArray+1] - d_u[id_numInArray-1])/(d_r[id_numInArray+1] - d_r[id_numInArray-1])) + (1-d_fourpoint_constant) * ((d_u[id_numInArray+2] - d_u[id_numInArray-2])/(d_r[id_numInArray+2] - d_r[id_numInArray-2]));
else
du = 0;
}
__syncthreads();
if((threadIdx.x > radius-1 && threadIdx.x < blockDim.x - radius) && (id_r < numNodesPerGrid) && (id_grid < numGrids))
{
d_u[ id_numInArray] = d_u[id_numInArray] * uchange; // if this line is commented out, results will be computed correctly ("Solution 2")
d_du[ id_numInArray] = du;
}
__syncthreads();
++id_grid;
}
}
bool gridValuesEqual(float *matarray, uint id0, uint id1, const char *label, int numNodesPerGrid){
bool retval = true;
for(uint i=0; idu_arr[k*paramH.numNodes + 668]:\t%E\n", du_arr[k*numNodesPerGrid + 668]);//*/
bool equal = true;
for(int k=1; k>>(...)\n\n", numBlocksPerSM * numSM, TPB, 0);
printf("Launching Kernel...\n\n");
heisenkernel<<>>(d_u, d_r, d_du, radius, numNodesPerGrid, numBlocksPerSM, numGridsPerSM, numGrids);
cudaDeviceSynchronize();
cudaMemcpy(u_arr, d_u, memsize_grid * numGrids, cudaMemcpyDeviceToHost);
cudaMemcpy(du_arr, d_du, memsize_grid * numGrids, cudaMemcpyDeviceToHost);
cudaMemcpy(r_arr, d_r, memsize_grid * numGrids, cudaMemcpyDeviceToHost);
/*
printf("\n\nafter kernel finished\n\n");
for(uint k=0; kdu_arr[k*paramH.numNodes + 668]:\t%E\n", du_arr[k*numNodesPerGrid + 668]);//*/
equal = true;
for(int k=1; k
Makefile文件:
CUDA = 1
DEFINES =
ifeq ($(CUDA), 1)
DEFINES += -DCUDA
CUDAPATH = /usr/local/cuda-6.5
CUDAINCPATH = -I$(CUDAPATH)/include
CUDAARCH = -arch=sm_20
endif
CXX = g++
CXXFLAGS = -pipe -g -std=c++0x -fPIE -O0 $(DEFINES)
VALGRIND = valgrind
VALGRIND_FLAGS = -v --leak-check=yes --log-file=out.memcheck
CUDAMEMCHECK = cuda-memcheck
CUDAMC_FLAGS = --tool memcheck
RACECHECK = $(CUDAMEMCHECK)
RACECHECK_FLAGS = --tool racecheck
INCPATH = -I. $(CUDAINCPATH)
LINK = g++
LFLAGS = -O0
LIBS =
ifeq ($(CUDA), 1)
NVCC = $(CUDAPATH)/bin/nvcc
LIBS += -L$(CUDAPATH)/lib64/
LIBS += -lcuda -lcudart -lcudadevrt
NVCCFLAGS = -g -G -O0 --ptxas-options=-v
NVCCFLAGS += -lcuda -lcudart -lcudadevrt -lineinfo --machine 64 -x cu $(CUDAARCH) $(DEFINES)
endif
all:
$(NVCC) $(NVCCFLAGS) $(INCPATH) -c -o $(DST_DIR)heisenbug.o $(SRC_DIR)heisenbug.cu
$(LINK) $(LFLAGS) -o heisenbug heisenbug.o $(LIBS)
clean:
rm heisenbug.o
rm heisenbug
memrace: all
./heisenbug > out
$(VALGRIND) $(VALGRIND_FLAGS) ./heisenbug > out.memcheck.log
$(CUDAMEMCHECK) $(CUDAMC_FLAGS) ./heisenbug > out.cudamemcheck
$(RACECHECK) $(RACECHECK_FLAGS) ./heisenbug > out.racecheck
Robert Crove..
9
请注意,在整篇文章中,我没有看到明确询问的问题,因此我回复:
我期待着了解这里发生了什么.
你有竞争条件d_u
.
通过你自己的声明:
•为了使块彼此独立,引入网格上的小重叠(每个网格的网格点666,667,668,669由来自不同块的两个线程读取,尽管只有一个线程正在写入他们,这是问题发生的重叠)
此外,如果您d_u
根据代码中的语句注释掉写入,则问题将消失.
CUDA线程块可以按任何顺序执行.您有至少2个不同的块从网格点666,667,668,669读取.结果将根据实际发生的情况而有所不同:
两个块在发生任何写入之前读取该值.
一个块读取该值,然后发生写入,然后另一个块读取该值.
如果一个块正在读取可由另一个块写入的值,则这些块不是彼此独立的(与您的语句相反).在这种情况下,块执行的顺序将决定结果,并且CUDA不指定块执行的顺序.
请注意,cuda-memcheck
该-tool racecheck
选项仅捕获与__shared__
内存使用相关的竞争条件.你发布的内核不使用任何__shared__
内存,因此我不希望cuda-memcheck
报告任何内容.
cuda-memcheck
为了收集数据,确实会影响块执行的顺序,因此它会影响行为并不奇怪.
in-kernel printf
代表一个代价高昂的函数调用,写入全局内存缓冲区.所以它也会影响执行行为/模式.如果要打印大量数据,超出输出的缓冲行,则在缓冲区溢出的情况下,效果非常高(就执行时间而言).
另外,据我所知,Linux Mint 不是CUDA支持的发行版.但是我不认为这与你的问题有关; 我可以在受支持的配置上重现该行为.
请注意,在整篇文章中,我没有看到明确询问的问题,因此我回复:
我期待着了解这里发生了什么.
你有竞争条件d_u
.
通过你自己的声明:
•为了使块彼此独立,引入网格上的小重叠(每个网格的网格点666,667,668,669由来自不同块的两个线程读取,尽管只有一个线程正在写入他们,这是问题发生的重叠)
此外,如果您d_u
根据代码中的语句注释掉写入,则问题将消失.
CUDA线程块可以按任何顺序执行.您有至少2个不同的块从网格点666,667,668,669读取.结果将根据实际发生的情况而有所不同:
两个块在发生任何写入之前读取该值.
一个块读取该值,然后发生写入,然后另一个块读取该值.
如果一个块正在读取可由另一个块写入的值,则这些块不是彼此独立的(与您的语句相反).在这种情况下,块执行的顺序将决定结果,并且CUDA不指定块执行的顺序.
请注意,cuda-memcheck
该-tool racecheck
选项仅捕获与__shared__
内存使用相关的竞争条件.你发布的内核不使用任何__shared__
内存,因此我不希望cuda-memcheck
报告任何内容.
cuda-memcheck
为了收集数据,确实会影响块执行的顺序,因此它会影响行为并不奇怪.
in-kernel printf
代表一个代价高昂的函数调用,写入全局内存缓冲区.所以它也会影响执行行为/模式.如果要打印大量数据,超出输出的缓冲行,则在缓冲区溢出的情况下,效果非常高(就执行时间而言).
另外,据我所知,Linux Mint 不是CUDA支持的发行版.但是我不认为这与你的问题有关; 我可以在受支持的配置上重现该行为.