OpenCL의 비용 get_local_id()

Question

나는 루프의 여러 블록 스캔을 계산하는 간단한 스캔 커널을 가지고있다. get_local_id()가 루프 내부에서 호출하는 대신 로컬 변수에 저장 될 때 성능이 다소 상승하는 것으로 나타났습니다. 그래서, 코드와이를 요약 :

__kernel void LocalScan_v0(__global const int *p_array, int n_array_size, __global int *p_scan) 
{ 
    const int n_group_offset = get_group_id(0) * SCAN_BLOCK_SIZE; 
    p_array += n_group_offset; 
    p_scan += n_group_offset; 
    // calculate group offset 

    const int li = get_local_id(0); // *** local id cached *** 
    const int gn = get_num_groups(0); 
    __local int p_workspace[SCAN_BLOCK_SIZE]; 
    for(int i = n_group_offset; i < n_array_size; i += SCAN_BLOCK_SIZE * gn) { 
     LocalScan_SingleBlock(p_array, p_scan, p_workspace, li); 

     p_array += SCAN_BLOCK_SIZE * gn; 
     p_scan += SCAN_BLOCK_SIZE * gn; 
    } 
    // process all the blocks in the array (each block size SCAN_BLOCK_SIZE) 
} 

는 GTX-780 74 GB/s의 처리량을 가지고,이 때 :

__kernel void LocalScan_v0(__global const int *p_array, int n_array_size, __global int *p_scan) 
{ 
    const int n_group_offset = get_group_id(0) * SCAN_BLOCK_SIZE; 
    p_array += n_group_offset; 
    p_scan += n_group_offset; 
    // calculate group offset 

    const int gn = get_num_groups(0); 
    __local int p_workspace[SCAN_BLOCK_SIZE]; 
    for(int i = n_group_offset; i < n_array_size; i += SCAN_BLOCK_SIZE * gn) { 
     LocalScan_SingleBlock(p_array, p_scan, p_workspace, get_local_id(0)); 
     // *** local id polled inside the loop *** 

     p_array += SCAN_BLOCK_SIZE * gn; 
     p_scan += SCAN_BLOCK_SIZE * gn; 
    } 
    // process all the blocks in the array (each block size SCAN_BLOCK_SIZE) 
} 

가 동일한 하드웨어에서 70 GB/s의를 가지고 있습니다. 유일한 차이점은 get_local_id()에 대한 호출이 루프의 내부인지 외부인지 여부입니다. LocalScan_SingleBlock()의 ​​코드는 this GPU Gems article에 설명되어 있습니다.

이제 몇 가지 질문이 제기됩니다. 나는 항상 스레드 ID가 일부 레지스터 내에 저장되어 있으며 스레드에 대한 액세스가 모든 스레드 로컬 변수만큼 빠르다는 것을 상상했습니다. 이것은 사실이 아닌 것 같습니다. 나는 항상 루프에서 함수를 호출하지 않을 오래된 "C"프로그래머가 꺼려하는 변수에 로컬 ID를 캐싱하는 습관을 가지고 있었지만 매번 동일한 값을 반환 할 것을 기대했지만, 심각하게 그것이 어떤 차이를 만들 것이라고 생각하지 않습니다.

이유에 대한 의견이 있으십니까? 나는 컴파일 된 바이너리 코드를 검사하지 않았다. 누구도 같은 경험을합니까? CUDA에서 threadIdx.x과 동일합니까? ATI 플랫폼은 어떻습니까? 이 동작이 어딘가에 설명되어 있습니까? 신속하게 CUDA Best Practices를 스캔했지만 아무것도 찾지 못했습니다.

Answer 1

이것은 단지 추측이지만 크로노스 페이지 당

http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/get_local_id.html

get_local_id()은 상수 값 (단순히이 size_t)을 반환하도록 정의되지 않는다. 즉, 컴파일러가 알고있는 한 컴파일러의 관점에서 볼 때 함수 값의 반환이 변경 될 수 있기 때문에 상수 local_id와 비교하여 특정 최적화를 수행하지 못하게 될 수 있습니다.)