我有以下算法:
__global__ void Update(int N, double* x, double* y, int* z, double* out)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < N)
{
x[i] += y[i];
if (y[i] >= 0.)
out[z[i]] += x[i];
else
out[z[i]] -= x[i];
}
}
重要的是要注意 out 小于 x。假设 x、y 和 z 总是相同的大小,比如 1000,out 总是更小,比如 100。z 是每个 x 和 y 对应的 in out 的索引。
除了out的更新外,这都是find。线程之间可能会发生冲突,因为 z 不仅包含唯一值而且有重复值。因此,我目前使用 atomicAdd 的原子版本实现了这一点。并使用比较和交换减去。这显然很昂贵,意味着我的内核运行时间要长 5-10 倍。
我想减少这个,但是我能想到的唯一方法是让每个线程都有自己的 out 版本(可以很大,10000+,X 10000+ 线程)。这意味着我设置 10000 double[10000](可能在共享中?)调用我的内核,然后对这些数组求和,也许在另一个内核中。肯定有更优雅的方法来做到这一点吧?
可能值得注意的是,x、y、z 和 out 驻留在全局内存中。由于我的内核(我有其他类似的内核)非常简单,我还没有决定跨位复制到共享(内核上的 nvvp 显示相等的计算和内存,所以我认为添加从移动数据的开销时不会获得太多性能全局到共享,然后再回来,有什么想法吗?)。
最佳答案
方法一:
构建一组“事务”。由于每个线程只有一个更新,因此您可以轻松地构建一个固定大小的“事务”记录,每个线程一个条目。假设我的
out
表中有 8 个线程(为简单起见)和一些任意数量的条目。假设我的 8 个线程想要执行 8 个这样的事务:thread ID (i): 0 1 2 3 5 6 7 z[i]: 2 3 4 4 3 2 3 x[i]: 1.5 0.5 1.0 0.5 0.1 -0.2 -0.1 "transaction": 2,1.5 3,0.5 4,1.0 4,0.5 3,0.1 2,-0.2 3,-0.1
现在对交易做一个 sort_by_key,按照
z[i]
的顺序排列它们:sorted: 2,1.5 2,-0.2 3,0.5 3,-0.1 3,0.1 4,1.0 4,0.5
现在对事务进行 reduce_by_key 操作:
keys: 2 3 4 values: 1.3 0.5 1.5
现在根据 key 更新
out[i]
:out[2] += 1.3 out[3] += 0.5 out[4] += 1.5
方法二:
如您所说,全局内存中有数组x
、y
、z
和out
。如果您打算重复使用作为“映射”的 z
,您可能需要按照 z
的顺序重新排列(分组)或排序数组:
index (i): 0 1 2 3 4 5 6 7
z[i]: 2 8 4 8 3 1 4 4
x[i]: 0.2 0.4 0.3 0.1 -0.1 -0.4 0.0 1.0
按 z[i] 分组:
index (i): 0 1 2 3 4 5 6 7
z[i]: 1 2 3 4 4 4 8 8
x[i]:-0.4 0.2 -0.1 0.3 0.0 1.0 0.4 0.1
这个,或者它的一些变体,可以让你不必在方法 1 中重复执行排序操作(同样,如果你重复使用相同的“映射” vector )。
关于c - 未知全局数组索引的归约或原子运算符,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/28555479/