并行化具有依赖性的代码

Question

我的代码如下：

数组a 包含有关其元素所属组的信息。即元素 i 属于组 a[i]。每个组可以包含两个元素。我想将此信息存储在长度为 2*number 个组的 b 中。因此，b[j] 和 b[j+1] 的值会给我属于组 j/2（整数除法）和 j 的元素是偶数。

void assign(int *a, int *b){
    for(int i = 0; i<N; i++){
        int group = a[i];
        int posit=2*group;
        if(b[2*i]!=0){
            posit++;
        }
        b[posit] = i;
    }
}

N as is clear length of a.

默认情况下，b[] 中的值为零，表示没有元素。

有明显的数据依赖性，并行化看起来并不容易。我正在寻找进一步的建议，如果我缺少一个聪明的方法。

Answer 1

通常，您可以尝试在包含对 {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 个元素，则可以重复此方法，但随着数量越来越多，它很快就不再可行了。

并行化具有依赖性的代码

Parallelizing a code with dependency

parallel-processing

cuda

hpc

openmp