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)
매개 변수
피 | 발 |
---|---|
셀 포인터 | 세포에 추가됨 |
비고
성능은 원자 연산 수 및 메모리 공간에 따라 다릅니다. 직렬 작업을하는 것은 gpu가 SIMD 배열이기 때문에 거의 항상 커널 실행 속도를 늦추고 배열의 각 유닛은 동일한 유형의 작업을 수행하지 않으면 다른 유닛을 대기합니다.
원자 추가 기능
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 구현에 따라 값을 임의 또는 전체적으로 순서대로 배치 할 수 있습니다.
Modified text is an extract of the original Stack Overflow Documentation
아래 라이선스 CC BY-SA 3.0
와 제휴하지 않음 Stack Overflow