opencl
원자력 운영

통사론

• int atomic_add (휘발성 __global int * p, int val)

• 부호없는 int atomic_add (휘발성 __global unsigned int * p, 부호없는 int val)

• int atomic_add (휘발성 __local int * p, int val)

• 부호없는 int atomic_add (휘발성 __local unsigned int * p, 부호없는 int val)

매개 변수

비고

원자 추가 기능

    int fakeMalloc(__local int * addrCounter,int size)
    {
        // lock addrCounter
        // adds size to addrCounter's pointed cell value
        // unlock
        // return old value of addrCounter's pointed cell

        // serial between all threads visiting -> slow
        return atomic_add(addrCounter,size);
    }

    __kernel void vecAdd(__global float* results )
    {
       int id = get_global_id(0);
       int lid=get_local_id(0);
       __local float stack[1024];
       __local int ctr;
       if(lid==0)
          ctr=0;
       barrier(CLK_LOCAL_MEM_FENCE);
       stack[lid]=0.0f;                    // parallel operation
       barrier(CLK_LOCAL_MEM_FENCE);
       int fakePointer=fakeMalloc(&ctr,1); // serial operation
       barrier(CLK_LOCAL_MEM_FENCE);
       stack[fakePointer]=lid;             // parallel operation
       barrier(CLK_GLOBAL_MEM_FENCE);
       results[id]=stack[lid];
    }

첫 번째 요소의 출력 :

때때로

192 193 194 195 196 197 198

때때로

0 1 2 3 4 5 6

때때로

128 129 130 131 132 133 134

지역 범위 = 256 인 설정입니다.

어떤 스레드가 먼저 fakeMalloc을 방문하든간에 첫 번째 결과 셀에 자체 로컬 스레드 ID를 저장합니다. 예제 gpu의 내부 SIMD 및 웨이브 프론트 구조는 인접한 64 개의 스레드가 인접한 결과 셀에 값을 배치 할 수있게합니다. 다른 장치는 해당 장치의 opencl 구현에 따라 값을 임의 또는 전체적으로 순서대로 배치 할 수 있습니다.