実行を高速化するために、ローカルメモリを使用してopenclでカーネルを作成しました。ローカルメモリを使用するのはこれが初めてです。私のglobal_work_size=16およびlocal_work_size=8。
Openclカーネル:mapper.cl
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
typedef struct data
{
double dattr[10];
int d_id;
int bestCent;
}Data;
typedef struct cent
{
double cattr[5];
int c_id;
}Cent;
__kernel void mapper(__global int *keyMobj, __global int *valueMobj, __global Data *dataMobj, __global Cent *centMobj)
{
int bx = get_group_id(0);
int tx = get_local_id(0);
int size = get_local_size(0);
__local double localData[8][2];
__local double localCent[2][2];
__local int local_id[8];
int index = tx + bx*size;
int j,k,color=0;
double dmin=1000000.0, dx;
for(j=0; j<2; j++)
{
if(tx<2)
localCent[tx][j] = centMobj[tx].cattr[j];
localData[tx][j] = dataMobj[index].dattr[j];
local_id[tx] = dataMobj[index].d_id;
}
barrier(CLK_LOCAL_MEM_FENCE);
for(j=0; j<2; j++)
{
dx = 0.0;
for(k=0; k<2; k++)
dx+= ((localCent[j][k] - localData[index][k]) * (localCent[j][k] - localData[index][k]));
if(dx<dmin)
{ color = j;
dmin = dx;
}
}
keyMobj[index] = color;
valueMobj[index] = local_id[tx];
}
上記のカーネルでは、構造体dataMobjの最初の8つのオブジェクトとcentMobjの2つのオブジェクトを、それぞれグローバルメモリからローカルメモリlocalDataとlocalCentにフェッチしています。私のプログラミングに関する限り、データフェッチ用の適切なforループを使用しました。
ただし、最初の8つの作業項目(最初の作業グループの場合)のみが正しく実行されます。残りのブロックの次の残りの8つの作業項目を使用してデータをフェッチしません。
これが正確にどのように機能するかを教えてください。問題が発生した場合はお知らせください。また、スレッドの同期にメモリフェンスを使用しました。