哪种方式订购共享的 2D/3D 阵列，以便在 CUDA/OpenCL 中对 1 维进行并行缩减

Question

总体目标

我要对二分图进行几次缩减，由两个用于顶点的密集数组和一个指定是否存在边的密集数组表示。比如说，两个数组是a0[]和a1[]，所有边都像e[i0][i1]一样(即从a0中的元素到a1中的元素(。

有 ~100+100 个

顶点和 ~100*100 条边，因此每个线程负责一条边。

任务1：最大缩减

对于a0中的每个顶点，我想找到连接到它的所有顶点(以a1为单位(的最大值，然后反过来也是如此：将结果分配给数组b0，对于a1中的每个顶点，我想找到连接顶点的最大b0[i0]。

为此，我：

1. 加载到共享内存中

    #define DC_NUM_FROM_SHARED 16
    #define DC_NUM_TO_SHARED 16
    __global__ void max_reduce_down(
            Value* value1
        , Value* max_value_in_connected
        , int r0_size, int r1_size
        , bool** connected
        )
    {
        int id_from;
        id_from = blockIdx.x * blockDim.x + threadIdx.x;
        id_to   = blockIdx.y * blockDim.y + threadIdx.y;
        bool within_bounds = (id_from < r0_size) && (id_to < r1_size);
        //load into shared memory
        __shared__ Value value[DC_NUM_TO_SHARED][DC_NUM_FROM_SHARED]; //FROM is the inner (consecutive) dimension
        if(within_bounds)
            value[threadIdx.y][threadIdx.x] = connected[id_to][id_from]? value1[id_to] : 0;
        else
            value[threadIdx.y][threadIdx.x] = 0;
        __syncthreads();
        if(!within_bounds)
            return;
   

2. 减少

    for(int stride = DC_NUM_TO_SHARED/2; threadIdx.y < stride; stride >>= 1)
    {
        value[threadIdx.y][threadIdx.x] = max(value[threadIdx.y][threadIdx.x], dc[threadIdx.y + stride][threadIdx.x]);
        __syncthreads();
    }
   

3. 回信

    max_value_connected[id_from] = value[0][threadIdx.x];
   

任务2：最佳k

类似的问题，但归约仅适用于a0中的顶点，我需要找到从连接a1中选择k最佳候选者(k为 ~5(。

1. 我用除第一名以外的零元素初始化共享数组

    int id_from, id_to;
    id_from = blockIdx.x * blockDim.x + threadIdx.x;
    id_to   = blockIdx.y * blockDim.y + threadIdx.y;
    __shared Value* values[MAX_CHAMPS * CHAMPS_NUM_FROM_SHARED * CHAMPS_NUM_TO_SHARED]; //champion overlaps
    __shared int* champs[MAX_CHAMPS * CHAMPS_NUM_FROM_SHARED * CHAMPS_NUM_TO_SHARED]; // overlap champions
   
    bool within_bounds = (id_from < r0_size) && (id_to < r1_size);
    int i = threadIdx.y * CHAMPS_NUM_FROM_SHARED + threadIdx.x;
    if(within_bounds)
    {
        values[i] = connected[id_to][id_from] * values1[id_to];
        champs[i] = connected[id_to][id_from] ? id_to : -1;
    }
    else
    {
        values[i] = 0;
        champs[i] = -1;
    }
    for(int place = 1; place < CHAMP_COUNT; place++)
    {
        i = (place * CHAMPS_NUM_TO_SHARED + threadIdx.y) * CHAMPS_NUM_FROM_SHARED + threadIdx.x;
        values[i] = 0;
        champs[i] = -1;
    }
    if(! within_bounds)
        return;
    __syncthreads();
   

2. 减少它

    for(int stride = CHAMPS_NUM_TO_SHARED/2; threadIdx.y < stride; stride >>= 1)
    {
        merge_2_champs(values, champs, CHAMP_COUNT, id_from, id_to, id_to + stride);
        __syncthreads();
    }
   

3. 写回结果

    for(int place = 0; place < LOCAL_DESIRED_ACTIVITY; place++)
        champs0[place][id_from] = champs[place * CHAMPS_NUM_TO_SHARED * CHAMPS_NUM_FROM_SHARED + threadIdx.x];
   

问题

• 如何对共享数组中的元素进行排序(转置(，以便内存访问更好地使用缓存？
• 在这一点上是否重要，或者我可以从其他优化中获得更多？
• 如果我需要针对任务 2 进行优化，转置边缘矩阵会更好吗？据我了解，任务 1 中存在对称性，因此在那里无关紧要。

附言

我在加载时延迟了展开循环并进行第一次缩减迭代，因为在我探索更简单的方法之前，我认为这太复杂了。

对于任务 2，最好不要加载零元素，因为数组永远不需要增长，只有在完成 log( k( 步骤后才会开始收缩。这将使共享内存中的紧凑k倍！但我害怕由此产生的索引数学。

语法和正确性

不寻常的类型只是typedef int s/char s/等 - AFAIK，在 GPU 中，尽可能压缩它们是有意义的。我还没有运行代码，无需检查索引错误。

另外，我

正在使用 CUDA，但我也对 OpenCL 的观点感兴趣，因为我认为最好的解决方案应该是相同的，无论如何我将来都会使用 OpenCL。

Answer 1

好吧，我想我想通了。

我正在考虑的两种选择是在y维度上进行缩减，在x维度上独立，反之亦然(x维度是连续的维度(。在任何情况下，调度程序都能够沿 x 维度将线程组装成扭曲，因此可以保证一定的一致性。但是，让连贯性超越翘曲会很棒。此外，由于共享阵列的 2D/3D 性质，必须将维度限制为 16 甚至 8。

为了确保翘曲内的合并，调度程序必须沿 x 维组装翘曲。

如果减小超过 x 维，则在每次迭代后，翘曲中的活动线程数将减半。但是，如果在 y 维上减少，则活动翘曲的数量将减半。

所以，我需要减少 y 以上。

除非转置(负载(是最慢的，这是异常情况。

Answer 2

合并缓冲区读取非常重要;如果你不这样做，内核可能会慢 32 倍。如果这意味着能够执行重新排列通道，则值得执行它们(当然，重新排列通道也需要合并，但您通常可以利用共享本地内存来执行此操作(。