Fruchtermon と Reingold のレイアウト アルゴリズムの OpenCL 実装に取り組んでいます。既に実装した CPU バージョンと比較すると、ほとんど機能しています。repel
ただし、大きなグラフの場合、グラフ内の頂点のすべてのペア間の反発力を計算する関数にボトルネックがあることに気付きました。このボトルネックを軽減するために OpenCLfloat4
構造を使用することにしましたが、結果はまちまちです。
を使用できるようにコードを設定しましたがfloat4
、実際には (デバッグ用に) 1 つの位置のみを使用します。単一の位置 ( float_workers = 1
) のみを使用すると、関数は正しく動作します
ただし、設定するfloat_workers > 1
と、ますます奇妙な動作になります。
コードの何が問題なのかわかりません。カーネル全体が Java クラスの文字列に含まれているため、JOCL を使用してい gid *= " + String.valueOf(float_workers) + ";
ます。
行マーカー(1)から(2)は、正しい数のアイテム(float_workersに応じて1〜4アイテム)でfloat4を設定しています。このループは、GID をノードのペア (float_workers = 1 の場合) にマップするか、float_workers の方が高い場合は 2、3、または 4 つのノードのペアにマップします。ライン マーカー (2) から (3) は、実行されている実際の作業 (ノードをはじく量の計算) です。終了する行マーカー (3) は、結果を「変位」配列に設定しているため、後でノードの位置を更新できます。明らかに、カーネルをキューに入れるときにワーカーの数を正しく調整するので、それは問題ではないようです。
xPos[...]
それぞれが正確に 1 回しか設定できないと仮定するのは正しいですか? それがこれを壊すと私が考えることができる唯一のことです。他に何か見逃したことがありますか?
__kernel void repel(__global const float *optDist,
__global const int *nIndexes,
__global const float *xPos,
__global const float *yPos,
__global float *xDisp,
__global float *yDisp,
__global int *nodes, +
__global int *startIndexes,
__global int* totalWork){
int sub = nodes[0] - 1;
int work_dim = get_work_dim();
int gid = 0;
for(int i = 0; i < work_dim - 1; i++){
gid += get_global_id(i) * get_global_size(i);
}
gid += get_global_id(work_dim - 1);
gid *= " + String.valueOf(float_workers) + ";
int gid_start = gid;
if(gid >= totalWork[0]){return;}
int v,u,i;
v = u = i = 0;
float4 xDelta = (float4)(0.0f,0.0f,0.0f,0.0f);
float4 yDelta = (float4)(0.0f,0.0f,0.0f,0.0f);
int found = 0;
(1)for(i = 0; i < nodes[0]; i++){
if(found == " + String.valueOf(float_workers) + "){
break;
}
if(gid < startIndexes[i]){
v = i - 1;
u = (gid - startIndexes[i - 1]) + 1;
gid ++;
i--;
if(found == 0){
xDelta.s0 = xPos[v] - xPos[u];
yDelta.s0 = yPos[v] - yPos[u];
}
if(found == 1){
xDelta.s1 = xPos[v] - xPos[u];
yDelta.s1 = yPos[v] - yPos[u];
}
if(found == 2){
xDelta.s2 = xPos[v] - xPos[u];
yDelta.s2 = yPos[v] - yPos[u];
}
if(found == 3){
xDelta.s3 = xPos[v] - xPos[u];
yDelta.s3 = yPos[v] - yPos[u];
}
found++;
}
}
(2)
float4 deltaLength = sqrt((xDelta * xDelta) + (yDelta * yDelta));
float4 _optDist = (float4)(optDist[0], optDist[0], optDist[0], optDist[0]);
float4 force = _optDist / deltaLength;
float4 xResult = (xDelta / deltaLength) * force;
float4 yResult = (yDelta / deltaLength) * force;
(3)
if ((xDelta.s0 == 0) && (yDelta.s0 == 0)) {
xDisp[gid_start + 0] = 0;
yDisp[gid_start + 0] = 0;
}
else{
xDisp[gid_start + 0] = xResult.s0;
yDisp[gid_start + 0] = yResult.s0;
}
if(" + String.valueOf(float_workers) + " > 1){
if ((xDelta.s1 == 0) && (yDelta.s1 == 0)) {
xDisp[gid_start + 1] = 0;
yDisp[gid_start + 1] = 0;
}
else{
xDisp[gid_start + 1] = xResult.s1;
yDisp[gid_start + 1] = yResult.s1;
}
}
if(" + String.valueOf(float_workers) + " > 2){
if ((xDelta.s2 == 0) && (yDelta.s2 == 0)) {
xDisp[gid_start + 2] = 0;
yDisp[gid_start + 2] = 0;
}
else{
xDisp[gid_start + 2] = xResult.s2;
yDisp[gid_start + 2] = yResult.s2;
}
}
if(" + String.valueOf(float_workers) + " > 3){
if ((xDelta.s3 == 0) && (yDelta.s3 == 0)) {
xDisp[gid_start + 3] = 0;
yDisp[gid_start + 3] = 0;
}
else{
xDisp[gid_start + 3] = xResult.s3;
yDisp[gid_start + 3] = yResult.s3;
}
}
}