我最近发现了比赛检查的工具cuda内存检查,在 CUDA 5.0 中可用(cuda-memcheck --tool racecheck
,参见英伟达文档 http://docs.nvidia.com/cuda/cuda-memcheck/index.html#using-racecheck)。该工具可以检测 CUDA 内核中共享内存的竞争条件。
在调试模式下,该工具没有检测到任何东西,这显然是正常的。但是,在发布模式下(-O3
),根据问题的参数,我会收到错误。
下面是一个错误示例(第 22 行共享内存初始化,第 119 行赋值):
========= ERROR: Potential WAW hazard detected at __shared__ 0x0 in block (35, 0, 0) :
========= Write Thread (32, 0, 0) at 0x00000890 in ....h:119:void kernel_test3<float, unsigned int=4, unsigned int=32, unsigned int=64>(Data<float, unsigned int=4, unsigned int=32, unsigned int=64>*)
========= Write Thread (0, 0, 0) at 0x00000048 in ....h:22:void kernel_test3<float, unsigned int=4, unsigned int=32, unsigned int=64>(Data<float, unsigned int=4, unsigned int=32, unsigned int=64>*)
========= Current Value : 13, Incoming Value : 0
- 首先让我惊讶的是线程 ID。当我第一次遇到这个错误时,每个块包含 32 个线程(id 0 到 31)。那么为什么线程id 32会出现问题呢?我什至添加了额外的检查
threadIdx.x
,但这没有改变任何事情。
- 我使用共享内存作为临时缓冲区,每个线程处理自己的多维数组参数,例如
__shared__ float arr[SIZE_1][SIZE_2][NB_THREADS_PER_BLOCK]
。我真的不明白怎么可能存在任何竞争条件,因为每个线程都处理自己的共享内存部分。
- 将网格大小从 64 个块减少到 32 个块似乎可以解决该问题(每个块 32 个线程)。我不懂为什么。
为了了解发生了什么,我使用一些更简单的内核进行了测试。
让我向您展示一个产生此类错误的内核示例。基本上,这个内核使用SIZE_X*SIZE_Y*NTHREADS*sizeof(float)
字节的共享内存,每个 SM 可以使用 48KB 的共享内存。
test.cu
template <unsigned int NTHREADS>
__global__ void kernel_test()
{
const int SIZE_X = 4;
const int SIZE_Y = 4;
__shared__ float tmp[SIZE_X][SIZE_Y][NTHREADS];
for (unsigned int i = 0; i < SIZE_X; i++)
for (unsigned int j = 0; j < SIZE_Y; j++)
tmp[i][j][threadIdx.x] = threadIdx.x;
}
int main()
{
const unsigned int NTHREADS = 32;
//kernel_test<NTHREADS><<<32, NTHREADS>>>(); // ---> works fine
kernel_test<NTHREADS><<<64, NTHREADS>>>();
cudaDeviceSynchronize(); // ---> gives racecheck errors if NBLOCKS > 32
}
汇编:
nvcc test.cu --ptxas-options=-v -o test
如果我们运行内核:
cuda-memcheck --tool racecheck test
-
kernel_test<32><<<32, 32>>>();
:32 个块,32 个线程 => 不会导致任何明显的竞赛检查错误。
-
kernel_test<32><<<64, 32>>>();
:64 个块,32 个线程 => 导致 WAW 危险(threadId.x
= 32?!)和错误。
========= ERROR: Potential WAW hazard detected at __shared__ 0x6 in block (57, 0, 0) :
========= Write Thread (0, 0, 0) at 0x00000048 in ....h:403:void kernel_test(void)
========= Write Thread (1, 0, 0) at 0x00000048 in ....h:403:void kernel_test(void)
========= Current Value : 0, Incoming Value : 128
========= INFO:(Identical data being written) Potential WAW hazard detected at __shared__ 0x0 in block (47, 0, 0) :
========= Write Thread (32, 0, 0) at 0x00000048 in ....h:403:void kernel_test(void)
========= Write Thread (0, 0, 0) at 0x00000048 in ....h:403:void kernel_test(void)
========= Current Value : 0, Incoming Value : 0
那么我在这里缺少什么?我在共享内存方面做错了什么吗? (我还是个初学者)
Update:
问题似乎来自于cudaDeviceSynchronize()
when NBLOCKS > 32
。为什么会发生这种情况?