로컬 메모리에 대한 액세스/동기화

Question

저는 GPGPU 프로그래밍을 처음 접했습니다. 동기화를 많이 필요로하는 알고리즘을 구현하려고합니다. 하나의 작업 그룹 (글로벌 및 로컬 크기가 동일한 값을 가짐)을 사용합니다.

문제가 생깁니다 : 문제점의 크기가 초과 될 때까지 프로그램이 올바르게 작동합니다. addedValue [0]에 너무 [123 (123)와 같은 65에서 다른 값 (다음 diffrent 커널 33 instanction 행 값 및이 있는지 I는 (stimulationsOut의 검사)를 실현할 일부 조사이 시도 후

__kernel void assort(
__global float *array, 
__local float *currentOutput, 
__local float *stimulations, 
__local int *noOfValuesAdded, 
__local float *addedValue, 
__local float *positionToInsert, 
__local int *activatedIdx, 
__local float *range, 
int size, 
__global float *stimulationsOut 
) 
{ 
int id = get_local_id(0); 
if (id == 0) {...} 

barrier(CLK_LOCAL_MEM_FENCE); 

for (int i = 2; i < size; i++) 
{ 
    int maxIdx; 
    if (id == 0) 
    { 
    addedValue[0] = array[i]; 
    {...} 
    } 
    barrier(CLK_LOCAL_MEM_FENCE); 


    if (id < noOfValuesAdded[0]){...} 
    else 
     barrier(CLK_LOCAL_MEM_FENCE); 
    barrier(CLK_LOCAL_MEM_FENCE); 
    if (activatedIdx[0] == -2) {...} 
    else {...} 

    barrier(CLK_LOCAL_MEM_FENCE); 
    if (positionToInsert[0] != -1.0f) {...} 

    barrier(CLK_LOCAL_MEM_FENCE); 
    stimulationsOut[id] = addedValue[0]; 
    return; 
    } 

32, 123 ... 123 (33 번째 요소) 66 ... 66 66 66 66 66 .. (65 번째 요소) 127 ... .. 127 ...]

__gl obal float * array는 READ_ONLY이고, for 루프에서 처음으로 addedValue [0]을 변경하지 않습니다. 이 문제를 해결할 수있는 것은 무엇입니까?

내 GPU 사양 : [https://devtalk.nvidia.com/default/topic/521502/gt650m-a-kepler-part-/]

두 주석 이후의 경우 몸의 문제는 reccuring되지 않습니다

  /*if (activatedIdx[0] == -2) 
     { 
      if (noOfValuesAdded[0] == 2) 
      { 
       positionToInsert[0] = 0.99f; 
      } 
      else if (id != 0 && id != maxIdx 
        && stimulations[id] >= stimulations[(id - 1)] 
        && stimulations[id] >= stimulations[(id + 1)]) 
      { 
       if ((1.0f - (fabs((currentOutput[(id - 1)] - currentOutput[id]))/range[0])) < stimulations[(id - 1)]) 
        positionToInsert[0] = (float)id - 0.01f; 
        else 
       positionToInsert[0] = (float)id + 0.99f; 
      } 
     }*/ 

및

if (positionToInsert[0] != -1.0f) 
    { 
     float temp = 0.0f; 
     /*if ((float)id>positionToInsert[0]) 
     { 
      temp = currentOutput[id]; 
      barrier(CLK_LOCAL_MEM_FENCE); 
      currentOutput[id + 1] = temp; 
     } 
     else 
     { 
      barrier(CLK_LOCAL_MEM_FENCE); 
     }*/ 
     barrier(CLK_LOCAL_MEM_FENCE); 

     if (id == round(positionToInsert[0])) 
     { 
      currentOutput[id] = addedValue[0]; 
      noOfValuesAdded[0] = noOfValuesAdded[0] + 1; 
     } 
    } 

업데이트 : 을 장벽, 알고리즘 작품을 수정 한 후 적절하게 크기가 768을 초과 할 때까지 (이는 내 GPU에 2 배의 코어 수가 있음). 최대 작업 그룹 크기 인 1024 개의 요소까지 사용할 수 있습니다. 내가 놓친 게 있니?

Answer 1

워프의 모든 작업 항목은 잠금 단계에서 동일한 명령을 실행합니다. 엔비디아의 워프 크기는 32 개의 작업 항목입니다. 커널이 최대 32 개의 작업 항목을 올바르게 작동 시키면 장벽에 문제가 있음을 알 수 있습니다. barrier에 대한

문서 말 :

라도 장벽을 넘어 실행을 계속하도록 허용하기 전에이 기능을 실행해야하는 프로세서 에 커널을 실행하는 작업 그룹의 모든 작업 항목

.

이 문제는 커널에서 볼 수 있습니다. 여기 예를 들어 :

if ((float)id>positionToInsert[0]) 
{ 
    temp = currentOutput[id]; 
    barrier(CLK_LOCAL_MEM_FENCE); // <---- some work items may meet here 
    currentOutput[id + 1] = temp; 
} 
else 
{ 
    barrier(CLK_LOCAL_MEM_FENCE); // <---- other work items may meet here 
} 

당신은 아마이 문제를 해결할 수 : 크기가 768을 초과 할 때까지 장벽을 해결 한 후

if ((float)id>positionToInsert[0]) 
    temp = currentOutput[id]; 
barrier(CLK_LOCAL_MEM_FENCE); // <---- here all work items meet at the same barrier 
if ((float)id>positionToInsert[0]) 
    currentOutput[id + 1] = temp; 

Answer 2

, 알고리즘 (내 GPU에서 코어 이상하게 2 배 번호입니다) 제대로 작동합니다. 최대 작업 그룹 크기 인 1024 개의 요소까지 사용할 수 있습니다. 내가 놓친 게 있니?