Проблемы с ядром opencl float4

Question

Я работаю над реализацией OpenCL алгоритма компоновки Fruchtermon и Reingold, у меня он довольно хорошо работает по сравнению с версией процессора, которую я уже реализовал. Однако я заметил, что для больших графиков на repel функция - которая вычисляет силу отталкивания между каждой парой вершин в графе. Я решил попробовать использовать OpenCL float4 структура, чтобы уменьшить это узкое место, но я получаю смешанные результаты.

Я настроил код, чтобы я мог использовать float4в то время как на самом деле используется только 1 из позиций (для отладки). Когда я использую только одну позицию (float_workers = 1) функция работает правильно

Тем не менее, когда я установил float_workers > 1 Я получаю все более странное поведение.

Я не могу понять, что не так с моим кодом - я использую JOCL, так как все ядро ​​находится в строке в моем классе Java, поэтому такие строки: gid *= " + String.valueOf(float_workers) + "; имеет смысл и компилировать.

линейные маркеры (1) - (2) устанавливают float4 с правильным количеством элементов (1-4 элемента, в зависимости от float_workers). Этот цикл отображает GID на пару узлов (когда float_workers = 1) или на две, три или четыре пары узлов, когда float_workers выше. Линейные маркеры (2) - (3) являются фактической выполняемой работой (вычисляя количество, чтобы оттолкнуть узлы). Линейные маркеры (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;
               }
             }
    }