Случай переключения CUDA через blockIdx - PullRequest
Новая
       5

Случай переключения CUDA через blockIdx

0 голосов
/ 14 октября 2019

У меня есть ядро ​​в форме: 

\\ x, y, z = ... \\ read from global memory

double res;
switch(blockIdx.y){
    case 0:
        res = x * y;
        break;
    case 1:
        res = y * z;
        break;
    ...
}

\\ ... = res \\ write to global memory

Количество дел достигает ~ 50, все помечены как последовательные (!) Целые числа, начиная с 0.

Когда я делаю .ptxэто показывает, что switch-case расширен в условные ветви, и это чувствуется. Дело в том, что даже если расхождение ветвей никогда не происходит (и не может произойти из-за конструкции), само число сравнений ухудшает производительность.

Чего я действительно хотел бы добиться, так это равномерных прыжков. Один из способов, который я вижу, состоит в том, чтобы отделить тело каждого случая в функции, создать массив функциональных указателей, а затем (в ядре) выбрать с помощью blockIdx.y. С другой стороны, функции очень малы (самое большее горстка умножений), поэтому были бы заметны накладные расходы на сам вызов функции.

Есть ли способ обойти это?

Редактировать.

Выбирается для определения каждого случая как отдельной функции устройства: 

__device__ double func0(const double x, const double y, const double z) { return x * y; }

и размещение массива функциональных указателей в постоянной памяти: 

__constant__ double (*const fptr[]) (const double, const double, const double) = { func0, func1, ... }

Чтобы при вызове ядра выбор и применение функции сливались с записью в глобальную память: 

... = fptr[blockIdx.y](x, y, z);

Что приводит кв гораздо более чистом коде и, по-видимому, без условного ветвления. Однако не наблюдалось ни прироста производительности, ни ухудшения, что заставляет меня задуматься, существенно ли отличается код, выполняемый на GPU в конце (с заменой blockIdx.y) для этих двух подходов.

...