我正在OpenCL中編寫一個程序,它接收兩個點的數組,並計算每個點的最近鄰居。OpenCL內存優化 - 最近的鄰居
我有兩個這樣的程序。其中一人將計算4維的距離,6維的距離。他們在下面:
4個維度:
kernel void BruteForce(
global read_only float4* m,
global float4* y,
global write_only ushort* i,
read_only uint mx)
{
int index = get_global_id(0);
float4 curY = y[index];
float minDist = MAXFLOAT;
ushort minIdx = -1;
int x = 0;
int mmx = mx;
for(x = 0; x < mmx; x++)
{
float dist = fast_distance(curY, m[x]);
if (dist < minDist)
{
minDist = dist;
minIdx = x;
}
}
i[index] = minIdx;
y[index] = minDist;
}
6個維度:
kernel void BruteForce(
global read_only float8* m,
global float8* y,
global write_only ushort* i,
read_only uint mx)
{
int index = get_global_id(0);
float8 curY = y[index];
float minDist = MAXFLOAT;
ushort minIdx = -1;
int x = 0;
int mmx = mx;
for(x = 0; x < mmx; x++)
{
float8 mx = m[x];
float d0 = mx.s0 - curY.s0;
float d1 = mx.s1 - curY.s1;
float d2 = mx.s2 - curY.s2;
float d3 = mx.s3 - curY.s3;
float d4 = mx.s4 - curY.s4;
float d5 = mx.s5 - curY.s5;
float dist = sqrt(d0 * d0 + d1 * d1 + d2 * d2 + d3 * d3 + d4 * d4 + d5 * d5);
if (dist < minDist)
{
minDist = dist;
minIdx = index;
}
}
i[index] = minIdx;
y[index] = minDist;
}
我正在尋找方法來優化這個程序GPGPU。我已經閱讀了一些關於GPGPU優化的文章(包括http://www.macresearch.org/opencl_episode6,附帶源代碼),通過使用本地內存。我試着將它與這個代碼想出了:
kernel void BruteForce(
global read_only float4* m,
global float4* y,
global write_only ushort* i,
__local float4 * shared)
{
int index = get_global_id(0);
int lsize = get_local_size(0);
int lid = get_local_id(0);
float4 curY = y[index];
float minDist = MAXFLOAT;
ushort minIdx = 64000;
int x = 0;
for(x = 0; x < {0}; x += lsize)
{
if((x+lsize) > {0})
lsize = {0} - x;
if ((x + lid) < {0})
{
shared[lid] = m[x + lid];
}
barrier(CLK_LOCAL_MEM_FENCE);
for (int x1 = 0; x1 < lsize; x1++)
{
float dist = distance(curY, shared[x1]);
if (dist < minDist)
{
minDist = dist;
minIdx = x + x1;
}
}
barrier(CLK_LOCAL_MEM_FENCE);
}
i[index] = minIdx;
y[index] = minDist;
}
我收到垃圾結果我的「我」輸出(例如許多值是一樣的)。任何人都可以指向正確的方向嗎?我將不勝感激任何有助於我改進此代碼的答案,或者可能會發現上述優化版本的問題。
非常感謝您 Cauê
你的數組通常有多大?對於6d版本,你介意最後兩個元素是完全使用還是完全刪除? – mfa
感謝您的回答!我的數組非常大(每個都有〜400000個點)。我不介意完全放棄最後兩個元素。我只是添加了它們,因爲在我的測試中,一個非對齊的版本(使用float *而不是float8 *)拋出「資源不足」錯誤 – Waneck
也是4D,我實際上只使用3個維度,所以這可能是也適用於此版本。 – Waneck