Вызовите несколько раз get_global_id () против сохранения результата в локальной переменной? - PullRequest
Новая
       5

Вызовите несколько раз get_global_id () против сохранения результата в локальной переменной?

7 голосов
/ 15 июля 2010

Это, наверное, глупый вопрос, но: Насколько дорого вызывать какую-то get_* функцию в OpenCL-ядрах? Лучше ли сохранить результат для будущего использования в каком-нибудь локальном файле или вызывать нужную функцию всякий раз, когда это необходимо?

Или это зависит от платформы?

PS Я думаю, что CUDA решает это лучше с различными переменными threadIdx.

1 Ответ

6 голосов
/ 16 июля 2010

Я думаю, что это должно быть бесплатно для всех архитектур GPU.Он должен быть заменен соответствующим аппаратным регистром или константой в банке кэша.

Компилятор также может выполнять постоянное распространение на нем.Вы можете проверить себя, используя AMD Stream Analyzer :

OpenCL: 

__kernel 
void testKernel(__global uint * uintArray)
{
    uint threadId = get_global_id(0);

    uintArray[threadId] = 0xbaadf00d;
}

Radeon HD 5870 (Cypress) в сборе: 

0 ALU: ADDR(32) CNT(10) KCACHE0(CB0:0-15) KCACHE1(CB1:0-15) 
      0  x: MOV         R1.x,  (0xBAADF00D, -0.001327039325f).x      
         t: MULLO_INT   ____,  R1.x,  KC0[1].x      
      1  x: ADD_INT     ____,  R0.x,  PS0      
      2  w: ADD_INT     ____,  PV1.x,  KC0[6].x      
      3  z: LSHL        ____,  PV2.w,  (0x00000002, 2.802596929e-45f).x      
      4  y: ADD_INT     ____,  KC1[0].x,  PV3.z      
      5  x: LSHR        R0.x,  PV4.y,  (0x00000002, 2.802596929e-45f).x      
01 MEM_RAT_CACHELESS_STORE_RAW: RAT(1)[R0].x___, R1,  VPM

Здесьget_global_id(0) соответствует постоянному значению банка кэша KC0[1].x.Поэтому для ответа на ваш вопрос я бы использовал наиболее читаемую форму.

...