Есть ли способ разворачивать петли в ядре AMD OpenCL с помощью компилятора?

Question

Я пытаюсь оценить различия в производительности между OpenCL для AMD и графических процессоров Nvidia. У меня есть ядро, которое выполняет умножение матрицы-вектора. Я запускаю ядро ​​на двух разных системах в тот момент, мой ноутбук с NVidia GT525m с Ubuntu 12.04 и CUDA 4.0 (который содержит библиотеки и заголовки OpenCL), а другой - настольный компьютер с AMD Radeon HD7970 снова с Ubuntu 12.04 и новейшие драйверы Catalyst.

В ядре у меня есть два оператора #pragma unroll, которые производят большое ускорение для реализации Nvidia OpenCL (~ 6x). Однако версия AMD OpenCL не дает ускорения. Глядя на ядро ​​с помощью анализатора ядра AMD APP, вы получаете сообщение об ошибке, что разворот не используется, потому что количество поездок неизвестно. Так что мой вопрос: #pragma unroll работает с AMD OpenCL или есть альтернатива (возможно, флаг компилятора, о котором я не знаю). Я включил ядро ​​ниже

__kernel void mvKernel(__global float* a, const __global float* x, __global float* y, int m, int n) 
{ 
    float sum = 0.0f; 
    __global float* A; 
    int i; 
    int j = 0; 
    int indx = get_global_id(0); 
    __local float xs[12000]; 
#pragma unroll 
    for(i = get_local_id(0); i < n; i+= get_local_size(0)) { 
     xs[i] = x[i]; 
    } 
    barrier(CLK_LOCAL_MEM_FENCE); 
    A = &a[indx]; 
#pragma unroll 256 
    for(i = 0; i < n; i++) { 
     sum += xs[i] * A[j]; 
     j += m; 
    } 
    y[indx] = sum; 
} 

Это же ядро ​​производит правильные результаты в обеих реализациях, но команды #pragma раскатать ничего для AMD не делать (проверено, комментируя их).

Answer 1

Это не задокументировано, но оно должно фактически работать с #pragma unroll. Можете ли вы проверить журнал компилятора, чтобы узнать, применяется ли разворот? Я не уверен, использует ли анализатор ядра тот же самый компилятор, что и среда выполнения OpenCL, вы можете проверить.

В противном случае, если вы знаете, что n поставляется в блоках по 256, вы можете развернуть вручную, имея один цикл над блоками из 256 элементов и другой внутри с фиксированным размером 256, что может быть проще развернуть. Это, несомненно, решит проблему, что счетчик поездок не известен статически.

Однако, помните, что разворачивание цикла обычно не так велико для победы, так как у вас нет большого количества регистров для кэширования ваших вычислений. Повышенное давление в регистре от разворота петли может привести к разливу регистров, что еще более медленное. Вы должны проверить, насколько быстро ядро ​​находится на карте AMD. Более новый компилятор NVIDIA OpenCL также может не принести больше пользы из прагмы unroll.