cuda - 使用 OpenCL 的累积数组求和

Question

我正在使用 OpenCL 计算 n 维点之间的欧几里得距离。我得到两个 n 维点列表，我应该返回一个数组，其中只包含从第一个表中的每个点到第二个表中的每个点的距离。

我的方法是执行常规的双循环(对于 Table1 中的每个点{对于 Table2 中的每个点{...}}，然后对并行中的每对点进行计算。

然后欧几里德距离被分成 3 部分:
1.取点中各个维度的差异
2. 对差异进行平方(仍然适用于每个维度)
3. 将 2 中获得的所有值相加。
4. 取3中得到的值的平方根。(本例中省略了这一步。)

一切都像魅力一样，直到我尝试累积所有差异的总和(即，执行上述过程的第 3 步，下面代码的第 49 行)。

作为测试数据，我使用 DescriptorLists，每个 2 点:
DescriptorList1: 001,002,003,...,127,128; (p1)
129,130​​,131,...,255,256； (p2)

DescriptorList2: 000,001,002,...,126,127; (p1)
128,129,130​​,...,254,255； (p2)

所以结果向量应该有以下值:128, 2064512, 2130048, 128
现在我得到的随机数随着每次运行而变化。

我感谢任何帮助或指导我做错了什么。希望一切都清楚我正在工作的场景。

#define BLOCK_SIZE 128

typedef struct
{
    //How large each point is
    int length;
    //How many points in every list
    int num_elements;
    //Pointer to the elements of the descriptor (stored as a raw array)
    __global float *elements;
} DescriptorList;

__kernel void CompareDescriptors_deb(__global float *C, DescriptorList A, DescriptorList B, int elements, __local float As[BLOCK_SIZE])
{

    int gpidA = get_global_id(0);

    int featA = get_local_id(0);

    //temporary array  to store the difference between each dimension of 2 points
    float dif_acum[BLOCK_SIZE];

    //counter to track the iterations of the inner loop
    int loop = 0;

    //loop over all descriptors in A
    for (int i = 0; i < A.num_elements/BLOCK_SIZE; i++){

        //take the i-th descriptor. Returns a DescriptorList with just the i-th
        //descriptor in DescriptorList A
        DescriptorList tmpA = GetDescriptor(A, i);

        //copy the current descriptor to local memory.
        //returns one element of the only descriptor in DescriptorList tmpA
        //and index featA
        As[featA] = GetElement(tmpA, 0, featA);
        //wait for all the threads to finish copying before continuing
        barrier(CLK_LOCAL_MEM_FENCE);

        //loop over all the descriptors in B
        for (int k = 0; k < B.num_elements/BLOCK_SIZE; k++){
            //take the difference of both current points
            dif_acum[featA] = As[featA]-B.elements[k*BLOCK_SIZE + featA];
            //wait again
            barrier(CLK_LOCAL_MEM_FENCE);
            //square value of the difference in dif_acum and store in C
            //which is where the results should be stored at the end.
            C[loop] = 0;
            C[loop] += dif_acum[featA]*dif_acum[featA];
            loop += 1;
            barrier(CLK_LOCAL_MEM_FENCE);
        }
    }
}

Answer 1

您的问题在于这些代码行:

C[loop] = 0;
C[loop] += dif_acum[featA]*dif_acum[featA];

您工作组中的所有线程(好吧，实际上是您的所有线程，但让我们稍后再谈)都试图在没有任何同步的情况下同时修改此数组位置。有几个因素使这真的很成问题:

工作组不能保证完全并行工作，这意味着对于某些线程 C[loop] = 0 可以在其他线程已经执行下一行

之后调用。

那些并行执行的人都从 C[loop] 读取相同的值，用它们的增量修改它并尝试写回相同的地址。我不完全确定写回的结果是什么(我认为其中一个线程成功写回，而其他线程失败，但我不完全确定)，但无论哪种方式都是错误的。

现在让我们解决这个问题:
虽然我们可以使用原子来让它在全局内存上工作，但它不会很快，所以让我们在本地内存中积累:

local float* accum;
...
accum[featA] = dif_acum[featA]*dif_acum[featA];
barrier(CLK_LOCAL_MEM_FENCE);
for(unsigned int i = 1; i < BLOCKSIZE; i *= 2)
{
    if ((featA % (2*i)) == 0)
        accum[featA] += accum[featA + i];
    barrier(CLK_LOCAL_MEM_FENCE);
}
if(featA == 0)
    C[loop] = accum[0];

当然，您可以为此重用其他本地缓冲区，但我认为这一点很明确(顺便说一句:您确定将在本地内存中创建 dif_acum 吗，因为我想我在某处读到了这不会放在本地内存中，这将使预加载 A 到本地内存有点毫无意义)。

关于此代码的其他一些要点:

您的代码似乎仅适用于工作组(您既不使用 groupid 也不使用 global id 来查看要处理的项目)，为了获得最佳性能，您可能希望使用更多。

可能是个人喜好，但对我来说，使用 get_local_size(0) 似乎更好对于工作组大小而不是使用定义(因为您可能会在主机代码中更改它而没有意识到您应该将您的 opencl 代码更改为)

代码中的屏障都是不必要的，因为没有线程访问本地内存中由另一个线程写入的元素。因此，您不需要为此使用本地内存。

考虑到最后一个子弹，你可以简单地做:

float As = GetElement(tmpA, 0, featA);
...
float dif_acum = As-B.elements[k*BLOCK_SIZE + featA];

这将使代码(不考虑前两个项目符号):

__kernel void CompareDescriptors_deb(__global float *C, DescriptorList A, DescriptorList B, int elements, __local float accum[BLOCK_SIZE])
{
   int gpidA = get_global_id(0);
   int featA = get_local_id(0);
   int loop = 0;
   for (int i = 0; i < A.num_elements/BLOCK_SIZE; i++){
       DescriptorList tmpA = GetDescriptor(A, i);
       float As = GetElement(tmpA, 0, featA);
       for (int k = 0; k < B.num_elements/BLOCK_SIZE; k++){
           float dif_acum = As-B.elements[k*BLOCK_SIZE + featA];

           accum[featA] = dif_acum[featA]*dif_acum[featA];
           barrier(CLK_LOCAL_MEM_FENCE);
           for(unsigned int i = 1; i < BLOCKSIZE; i *= 2)
           {
              if ((featA % (2*i)) == 0)
                 accum[featA] += accum[featA + i];
              barrier(CLK_LOCAL_MEM_FENCE);
           }
           if(featA == 0)
              C[loop] = accum[0];
           barrier(CLK_LOCAL_MEM_FENCE);

           loop += 1;
        }
    }
}