我正在使用 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);
}
}
}
最佳答案
您的问题在于这些代码行:
C[loop] = 0;
C[loop] += dif_acum[featA]*dif_acum[featA];
您工作组中的所有线程(好吧,实际上是您的所有线程,但让我们稍后再谈)都试图在没有任何同步的情况下同时修改此数组位置。有几个因素使这真的很成问题:
现在让我们解决这个问题:
虽然我们可以使用原子来让它在全局内存上工作,但它不会很快,所以让我们在本地内存中积累:
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 到本地内存有点毫无意义)。
关于此代码的其他一些要点:
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;
}
}
}
关于cuda - 使用 OpenCL 的累积数组求和,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/3770533/