我有一个 CUDA 程序，其内核基本上执行以下操作。
- 我提供了笛卡尔坐标中的 n 个点的列表，例如维度为 dim_x * dim_y 的平面中的 (x_i,y_i)。我相应地调用内核。
- 对于这个平面上的每个点 (x_p,y_p)，我通过一个公式计算这 n 个点中的每一个到达那里所需的时间；给定这 n 个点以一定的速度移动。
- 我以 t_0、t_1、...t_n 的递增顺序对这些时间进行排序，其中 t_i 的精度设置为 1。即如果 t'_i=2.3453，那么我将只使用 t_i=2.3。
- 假设时间是从正态分布生成的，我模拟了 3 个最快的时间，以找出这 3 个点最早到达的时间百分比。因此通过随机实验假设 prob_0 = 0.76,prob_1=0.20 和 prob_2=0.04。由于 t_0 在三者中排名第一，因此我还返回了该点的原始索引（在时间排序之前）。说 idx_0 = 5（整数）。
- 因此，对于该平面上的每个点，我都会得到一对（prob，idx）。
假设这些点中的 n/2 是一种，其余的是另一种。生成的示例图像如下所示。
特别是当时间精度设置为 1 时，我注意到唯一 3 个时间元组 (t_0,t_1,t_2) 的数量仅为总数据点的 2.5%，即平面上的点数。这意味着在大多数情况下，内核只能使用先前模拟的值进行无用的模拟。因此，我可以使用一个字典，其键为时间的三元组，值作为索引和概率。由于据我所知和测试，STL 无法在内核内部访问，因此我构造了一个大小为 201000000 的浮点数组。这个选择是通过实验得出的，因为前 3 次都没有超过 20 秒。因此 t_0 可以从 {0.0,0.1,0.2,...,20.0} 中取任何值，因此有 201 个选择。我可以为这样的字典构造一个键，如下所示
- 键 = t_o * 10^6 + t_1 * 10^3 + t_2
就值而言，我可以将其设为 (prob+idx)。由于 idx 是一个整数并且 0.0<=prob<=1.0，我可以稍后检索这两个值
- prob=dict[key]-floor(dict[key])
- idx = floor(dict[key])
所以现在我的内核如下所示
__global__ my_kernel(float* points,float* dict,float *p,float *i,size_t w,...){
unsigned int col = blockIdx.y*blockDim.y + threadIdx.y;
unsigned int row = blockIdx.x*blockDim.x + threadIdx.x;
//Calculate time taken for each of the points to reach a particular point on the plane
//Order the times in increasing order t_0,t_1,...,t_n
//Calculate Key = t_o * 10^6 + t_1 * 10^3 + t_2
if(dict[key]>0.0){
prob=dict[key]-floor(dict[key])
idx = floor(dict[key])
}
else{
//Simulate and find prob and idx
dict[key]=(prob+idx)
}
p[row*width+col]=prob;
i[row*width+col]=idx;
}
结果与大多数点的原始程序非常相似，但对于某些点来说是错误的。
我很确定这是由于比赛条件造成的。请注意，dict 被初始化为全零。基本思想是使数据结构在字典的特定位置“读取多次写入一次”。
我知道可能有更优化的方法来解决这个问题，而不是分配这么多的内存。在这种情况下请告诉我。但我真的很想了解为什么这个特定的解决方案会失败。特别是我想知道如何在这个设置中使用 atomicAdd 。我没有使用它。