algorithm - 在 CUDA/OpenCL 中，哪种方式可以订购共享 2D/3D 阵列以并行缩减 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 中的顶点，我需要找到 k 个最佳候选是从 a1 中的连接中选出的(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 的 ints/chars/etc - AFAIK，在 GPU 中，尽可能地压缩它们是有意义的。我还没有运行代码，不需要检查索引错误。

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

Answer 1

好的，我想我明白了。

我正在考虑的两个备选方案是对 y 维度进行缩减，并独立于 x 维度，反之亦然(x 维度是连续的)。在任何情况下，调度程序都能够沿 x 维度将线程组装成 warp，因此可以保证一定的一致性。然而，让连贯性延伸到扭曲之外会很棒。此外，由于共享数组的 2D/3D 性质，人们必须将维度限制为 16 甚至 8。

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

如果减少 x 维度，在每次迭代后，warp 中的事件线程数将减半。但是，如果减少 y 维度，则事件扭曲的数量将减半。

所以，我需要减少 y。

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