通常,您可以尝试在包含对{i, a[i]}
的数组上使用并行排序算法。也许有一种更快的通用方法,我目前没有看到......
但是,在CUDA中,您可以利用这一事实,即当您有2个冲突线程写入32位字相同的内存位置 - 一个保证成功(但你不知道哪一个)。形式上,它充当一台Arbitrary CRCW机器。
因此,你可以在2个内核调用解决您的问题:
__global__ void assign1(int* a, int* b, int elementCount) {
int idx = threadIdx.x + blockIdx.x*blockDim.x;
if (idx<elementCount) {
int group = a[idx];
b[2*group] = idx;
}
}
__global__ void assign2(int* a, int* b, int elementCount) {
int idx = threadIdx.x + blockIdx.x*blockDim.x;
if (idx<elementCount) {
int group = a[idx];
if (b[2*group] != idx)
b[2*group+1] = idx;
}
}
__host__ void assign(int* dev_a, int* dev_b, int elementCount) {
int gridSize = elementCount/512+1;
int blockSize = 512;
assign1<<<gridSize,blockSize>>>(dev_a, dev_b, elementCount);
assign2<<<gridSize,blockSize>>>(dev_a, dev_b, elementCount);
}
在assign1
最多2个线程写入同一个存储位置b[2*group]
。其中一个线程保证成功。在assign2
写入失败的线程重复使用b[2*group+1]
的过程。
如果组中最多有3个或4个元素,则可以重复使用此方法,但随着数量增加,可能会很快停止执行。