我有一个CUDA内核似乎有竞争条件，并试图找出这种竞争条件的来源。我知道cuda-memcheck的'racecheck'工具，但是racecheck告诉我使用小输入时没有危险，这实际上与我自己的调查一致。对于大输入，虽然racecheck似乎需要永远（字面意思），所以我不能使用它。简要说明，定义为d_mat_3d变量的1D向量__device__用0填充并加载到全局存储器中。作为内核输入的两个大数组（d_A和d_v）也在main中定义并传递给核心。切割一个名为d_mat_3d的数组mat_2d，加载到共享内存中，并对其进行一些处理。然后，mat_2d将在全球记忆中写回d_mat_3d。
如此处所示，原子操作被用作不使用原子操作，mat_2d将遇到不同线程的竞争条件。
我猜我还有某种竞争条件的原因是mat_3d的结果每次都不同。
有没有想过这种竞争条件可能来自哪里？我可以采取任何步骤来清除它（除了工具竞赛检查）？如果你认为，没有竞争条件的证据，你能解释为什么每次执行内核时为d_mat_3d分配不同的值吗？
CUDA 9.0 / NVidia Titan Black / Ubuntu 16.04
#include <cstdlib>
#include <sstream>
#include <cstdio>
#include <cuda.h>
#include <cuda_runtime_api.h>
#define W 7 // fix limit for loops in kernel
#define SIZE 100 // defining matrix dimension
#define N_ELEM 10000 // no of elements in each vector
#define NTPB 1024 // no of threads per block
using namespace std;
__device__ float d_mat_3d[SIZE*SIZE*SIZE];
__global__ void cuda_kernel(float *d_A, float *d_v){
__shared__ float mat_2d[SIZE*SIZE]; // a 2D slice of 3D matrix d_mat_3d
unsigned int n = blockDim.x*blockIdx.x+threadIdx.x;
if(n >= N_ELEM)
return;
int x, y, z, i;
float r;
float A = d_A[n];
float v = d_v[n];
#pragma unroll
for(x=0; x<SIZE; x++){
// load mat_2d (on shared memory) using d_mat_3d (on global memory)
for(i=0; i<SIZE*SIZE; i++){
mat_2d[i] = d_mat_3d[i+x*SIZE*SIZE];
}
// sync threads as mat_2d is on shared memory
__syncthreads();
for(y=SIZE/2; y<SIZE/2+W; y++){
for(z=SIZE/2; z<SIZE/2+W; z++){
r = sqrt( pow(A,2) / v ); // no need to be in these loops. I know, but for my real case, it must be.
atomicAdd(&mat_2d[z+y*SIZE], r); // atomically add r
}
}
__syncthreads();
// write mat_2d (shared memory) back to mat_3d (global memory)
for(i=0; i<SIZE*SIZE; i++){
d_mat_3d[i+x*SIZE*SIZE] = mat_2d[i];
}
}
}
// this function writes h_mat_3d to disk.
void write_image(float *h_mat_3d){
ostringstream o_addToFile;
o_addToFile << "mat3d.bin";
FILE *pFile;
pFile = fopen(o_addToFile.str().c_str(), "wb");
for(int i=0; i<SIZE*SIZE*SIZE; i++){
fwrite(&h_mat_3d[i], sizeof(float), 1, pFile);
}
fclose (pFile);
}
int main(){
int i;
float *h_A = new float[N_ELEM]; // some large vector
float *h_v = new float[N_ELEM]; // some other large vector
float h_mat_3d[SIZE*SIZE*SIZE]; // will be filled w/ 0
float *d_A; // device variables
float *d_v;
for(i=0; i<N_ELEM; i++){
h_A[i] = 0.2f+(float)i/N_ELEM; // fill out with some calculations
h_v[i] = 0.5f+2.f*i/N_ELEM;
}
for(i=0; i<SIZE*SIZE*SIZE; i++){
h_mat_3d[i] = 0.f; // fill h_mat_3d with 0
}
cudaMalloc((void **)&d_A, sizeof(float)*N_ELEM); // allocate variables on device
cudaMalloc((void **)&d_v, sizeof(float)*N_ELEM);
cudaMemcpy(d_A, h_A, sizeof(float)*N_ELEM, cudaMemcpyHostToDevice); // copy from host to device
cudaMemcpy(d_v, h_v, sizeof(float)*N_ELEM, cudaMemcpyHostToDevice);
cudaMemcpyToSymbol(d_mat_3d, &h_mat_3d, sizeof(float)*SIZE*SIZE*SIZE); // copy h_mat_3d to device
cuda_kernel<<<(N_ELEM+NTPB-1)/NTPB,NTPB>>>(d_A, d_v); // execute kernel
cudaMemcpyFromSymbol(h_mat_3d, d_mat_3d, sizeof(float)*SIZE*SIZE*SIZE); // write it back to h_mat_3d
write_image(h_mat_3d); // write h_mat_3d to disk for checking
cudaFree(d_A); // free memory
cudaFree(d_v);
delete [] h_A;
delete [] h_v;
return 0;
}
是的，您的代码中至少有2种不同的竞争条件。
__syncthreads()保护加载操作的开始和结束。这样做会减少从一次运行到第六次或第七次有效十进制数的变化，这与ordinary float variability in floating-point operations一致，其中操作顺序不重复（这通常是这里的情况）。
添加以下行：
for(x=0; x<SIZE; x++){
__syncthreads(); // add this line
// load mat_2d (on shared memory) using d_mat_3d (on global memory)
for(i=0; i<SIZE*SIZE; i++){
mat_2d[i] = d_mat_3d[i+x*SIZE*SIZE];
}
// sync threads as mat_2d is on shared memory
__syncthreads();
应该大多纠正这个问题。如果没有这个，当你的内核在x中循环时，一些warp可以“前进”以开始加载共享内存，而之前的warp仍然忙于x中的前一个循环迭代（并注意下面的注释2，这可能会加剧这个问题。）d_mat_3d，因此每个线程块都会尝试写入各种值，因此存在竞争条件。线程块执行的顺序（由CUDA未定义）将主要决定最终的结果，并且这可以很容易地改变运行。我知道在没有完整内核重写的情况下解决这个问题的唯一方法就是简单地启动1个threadblock（它仍将填充d_mat_3d的相同区域）。这种竞争条件是全球记忆竞赛，目前cuda-memcheck无法发现这种种族。我对此过程犹豫不决，但是这段代码没有任何意义，并且要么表示缺乏对合理代码的关注，要么缺乏对CUDA执行模型的理解（特别是结合下面的第2项）。 ）还有一些我要指出的其他事情。
__syncthreads()可能是非法的。这个结构：
if(n >= N_ELEM)
return;
将允许（最后一个）threadblock中的某些线程提前退出，这意味着它们不会参与后续的__syncthreads()语句。这在CUDA中是非法的，并且限制在the programming guide中涵盖。这可以通过删除早期返回来修复，并使用if (n < N_ELEM)或类似方法保护内核循环的各个段（__syncthreads（）语句除外）。我不是建议这涵盖代码的每个问题，只是我注意到的事情。这是一个相对完整的测试用例，我用它来验证我的发现。它包括一些更改，以解决我上面提到的项目，以及对我来说似乎很重要的各种其他更改：
$ cat t268.cu
#include <cstdlib>
#include <sstream>
#include <cstdio>
#include <cuda.h>
#include <cuda_runtime_api.h>
#define W 7 // fix limit for loops in kernel
#define SIZE 100 // defining matrix dimension
#define N_ELEM 10000 // no of elements in each vector
#define NTPB 1024 // no of threads per block
using namespace std;
__device__ float d_mat_3d[SIZE*SIZE*SIZE];
__global__ void cuda_kernel(float *d_A, float *d_v){
__shared__ float mat_2d[SIZE*SIZE]; // a 2D slice of 3D matrix d_mat_3d
unsigned int n = blockDim.x*blockIdx.x+threadIdx.x;
int x, y, z, i;
float r;
float A = d_A[n];
float v = d_v[n];
#pragma unroll
for(x=0; x<SIZE; x++){
__syncthreads();
if (n < N_ELEM){
// load mat_2d (on shared memory) using d_mat_3d (on global memory)
for(i=0; i<SIZE*SIZE; i++){
mat_2d[i] = d_mat_3d[i+x*SIZE*SIZE];
}
}
// sync threads as mat_2d is on shared memory
__syncthreads();
if (n < N_ELEM){
for(y=SIZE/2; y<SIZE/2+W; y++){
for(z=SIZE/2; z<SIZE/2+W; z++){
r = sqrt( pow(A,2) / v ); // no need to be in these loops. I know, but for my real case, it must be.
atomicAdd(&(mat_2d[z+y*SIZE]), r); // atomically add r
}
}
}
__syncthreads();
// write mat_2d (shared memory) back to mat_3d (global memory)
if (n < N_ELEM){
for(i=0; i<SIZE*SIZE; i++){
d_mat_3d[i+x*SIZE*SIZE] = mat_2d[i];
}
}
}
}
// this function writes h_mat_3d to disk.
void write_image(float *h_mat_3d){
for (int i = 0; i < SIZE*SIZE; i++){
for (int j = 0; j < SIZE; j++)
if (h_mat_3d[i*SIZE+j] > 1.0f) printf("%d:%f\n ", i*SIZE+j, h_mat_3d[i*SIZE+j]);
printf("\n");}
}
int main(){
int i;
float *h_A = new float[N_ELEM]; // some large vector
float *h_v = new float[N_ELEM]; // some other large vector
float *h_mat_3d = new float[SIZE*SIZE*SIZE]; // will be filled w/ 0
float *d_A; // device variables
float *d_v;
for(i=0; i<N_ELEM; i++){
h_A[i] = 0.2f+i/(float)N_ELEM; // fill out with some calculations
h_v[i] = 0.5f+2.f*i/(float)N_ELEM;
}
for(i=0; i<SIZE*SIZE*SIZE; i++){
h_mat_3d[i] = 0.f; // fill h_mat_3d with 0
}
cudaMalloc((void **)&d_A, sizeof(float)*N_ELEM); // allocate variables on device
cudaMalloc((void **)&d_v, sizeof(float)*N_ELEM);
cudaMemcpy(d_A, h_A, sizeof(float)*N_ELEM, cudaMemcpyHostToDevice); // copy from host to device
cudaMemcpy(d_v, h_v, sizeof(float)*N_ELEM, cudaMemcpyHostToDevice);
cudaMemcpyToSymbol(d_mat_3d, h_mat_3d, sizeof(float)*SIZE*SIZE*SIZE); // copy h_mat_3d to device
cuda_kernel<<<1,NTPB>>>(d_A, d_v); // execute kernel
cudaMemcpyFromSymbol(h_mat_3d, d_mat_3d, sizeof(float)*SIZE*SIZE*SIZE); // write it back to h_mat_3d
write_image(h_mat_3d); // write h_mat_3d to disk for checking
cudaFree(d_A); // free memory
delete [] h_A;
delete [] h_v;
return 0;
}
$ nvcc -arch=sm_52 -o t268 t268.cu
$ ./t268 > out1.txt
$ ./t268 > out2.txt
$ diff out1.txt out2.txt |more
51,57c51,57
< 5050:330.657715
< 5051:330.657715
< 5052:330.657715
< 5053:330.657715
< 5054:330.657715
< 5055:330.657715
< 5056:330.657715
---
> 5050:330.657654
> 5051:330.657593
> 5052:330.657593
> 5053:330.657593
> 5054:330.657593
> 5055:330.657593
> 5056:330.657593
59,65c59,65
< 5150:330.657715
< 5151:330.657715
< 5152:330.657715
< 5153:330.657715
< 5154:330.657745
< 5155:330.657745
< 5156:330.657745
---
> 5150:330.657593
> 5151:330.657593
> 5152:330.657593
> 5153:330.657593
> 5154:330.657593
> 5155:330.657593
> 5156:330.657593
67,73c67,73
< 5250:330.657745
< 5251:330.657745
< 5252:330.657745
< 5253:330.657745
< 5254:330.657715
< 5255:330.657715
< 5256:330.657715
---
> 5250:330.657593
> 5251:330.657593
> 5252:330.657623
> 5253:330.657593
> 5254:330.657593
> 5255:330.657593
> 5256:330.657593
75,81c75,81
< 5350:330.657715
< 5351:330.657715
< 5352:330.657715
< 5353:330.657715
< 5354:330.657715
< 5355:330.657745
< 5356:330.657715
---
> 5350:330.657593
> 5351:330.657593
$
可以看出，剩余的变化是第7个有效十进制数字：
51,57c51,57
< 5050:330.657715
...
---
> 5050:330.657654