我有一个在两个层面上平行的问题：我有大量的(x0, x1, y0, y1)坐标对集合，它们被转换为变量vdx, vdy, vyy，对于这些集合中的每一个，我都试图计算由它们组成的所有“单项式”的值n度（即它们的不同幂的所有可能组合，例如vdx^3*vdy*vyy^2或vdx*1*vyy^4）。然后将这些值累加到所有集合中。
我的策略（现在我只想让它工作，它不必使用多个内核或复杂的缩减进行优化，除非它真的必须）是让每个线程处理一组坐标对并计算所有对应的单项式的值。每个块的共享内存保存所有单项式和，当块完成时，块中的第一个线程将结果添加到全局和。由于所有地方的所有线程都可以访问每个块的共享内存，因此我使用的是atomicAdd; 与块和全局内存相同。
不幸的是，某处似乎仍然存在竞争条件，因为每次运行内核时我都会得到不同的结果。
如果有帮助，我目前正在使用degree = 3并省略其中一个变量，这意味着在下面的代码中，最里面的 for 循环（over evbl）不做任何事情，只是重复 4 次。确实，内核的输出看起来像这样：51502,55043.1,55043.1,51502,47868.5,47868.5,48440.5,48440.6,46284.7,46284.7,46284.7,46284.7,46034.3,46034.3,46034.3,46034.3,44972.8,44972.8,44972.8,44972.8,43607.6,43607.6,43607.6,43607.6,43011,43011,43011,43011,42747.8,42747.8,42747.8,42747.8,45937.8,45937.8,46509.9,46509.9,...值得注意的是，有一个（粗略的）4 元组模式。但是每次我运行它时，值都非常不同。
一切都在浮动，但我在一个 2.1 GPU 上，所以这不应该是一个问题。cuda-memcheck 也没有报告错误。
有更多 CUDA 经验的人可以给我一些指示如何在这里追踪比赛条件吗？
__global__ void kernel(...) {
extern __shared__ float s_data[];
// just use global memory for now
// get threadID:
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if(idx >= nPairs) return;
// ... do some calculations to get x/y...
// calculate vdx, vdy and vyy
float vdx = (x1 - x0)/(float)xheight;
float vdy = (y1 - y0)/(float)xheight;
float vyy = 0.5*(y0 + y1)/(float)xheight;
const int offs1 = degree + 1;
const int offs2 = offs1 * offs1;
const int offs3 = offs2 * offs1;
float sol = 1.0;
// now calculate monomial results and store in shared memory
for(int evdx = 0; evdx <= degree; evdx++) {
for(int evdy = 0; evdy <= degree; evdy++) {
for(int evyy = 0; evyy <= degree; evyy++) {
for(int evbl = 0; evbl <= degree; evbl++) {
s = powf(vdx, evdx) + powf(vdy, evdy) + powf(vyy, evyy);
atomicAdd(&(s_data[evbl + offs1*evyy + offs2*evdy +
offs3*evdx]), sol/1000.0 );
}
}
}
}
// now copy shared memory to global
__syncthreads();
if(threadIdx.x == 0) {
for(int i = 0; i < nMonomials; i++) {
atomicAdd(&outmD[i], s_data[i]);
}
}
}