我正在研究 Fruchtermon 和 Reingold 布局算法的 OpenCL 实现,与我已经实现的 CPU 版本相比,我的工作非常好。但是我注意到,对于大型图,repel
函数存在瓶颈 - 它计算图中每对顶点之间的排斥力。我决定尝试使用 OpenCLfloat4
结构来减少这个瓶颈,但结果好坏参半。
我已经设置了代码,所以我可以使用float4
,而实际上只使用了 1 个位置(用于调试)。当我仅使用一个位置 ( float_workers = 1
) 时,该功能正常工作
但是,当我设置时,我float_workers > 1
的行为变得越来越奇怪。
我无法弄清楚我的代码有什么问题 - 我正在使用 JOCL,因为整个内核都在我的 Java 类中的一个字符串中,所以诸如: gid *= " + String.valueOf(float_workers) + ";
做有意义,并编译之类的行。
线标记 (1) 到 (2) 正在使用正确数量的项目(1-4 个项目,取决于 float_workers)设置 float4。此循环将 GID 映射到一对节点(当 float_workers = 1 时)或当 float_workers 较高时映射到两对、三对或四对节点。线标记 (2) 到 (3) 是正在执行的实际工作(计算排斥节点的量)。到 end 的线标记 (3) 将结果设置为“位移”数组,因此可以稍后更新节点的位置。显然,当我将内核排入队列时,我会正确调整工人的数量,所以这似乎不是问题。
我是否正确假设每个xPos[...]
只能设置一次?这是我能想到的唯一能打破这一点的事情。我错过了什么吗?
__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;
}
}
}