Ядро CUDA для оптимизации

Question

Привет, у меня недавно было ядро ​​CUDA для оптимизации. Вот оригинал CUDA ядра:

__glboal__ void kernel_base(float *data, int x_dim, int y_dim) 
{ 
    int ix = blockIdx.x; 
    int iy = blockIdx.y*blockDim.y + threadIdx.y; 
    int idx = iy*x_dim + ix; 
    float tmp = data[idx]; 

    if(ix % 2) 
    { 
    tmp += sqrtf(sinf(tmp) + 1.f); 
    } 
    else 
    { 
    tmp += sqrtf(cosf(tmp) + 1.f); 
    } 
    data[idx] = tmp; 
} 


dim3 block(1, 512); 
dim3 grid(2048/1, 2048/512); 
kernel<<<grid,block>>>(d_data, 2048, 2048); 

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

У кого-нибудь есть идея, как повысить производительность?

Answer 1

Примите во внимание, что

sin(x) = cos(x + pi/2) 

Соответственно, вы можете заменить if ... else условия

tmp += sqrtf(cosf(tmp + (ix%2) * pi/2) + 1.f); 

избежать ветви дивергенции.

Answer 2

Если бы я делал это, я бы сделал размеры блоков 16 x 16 или другую форму с более низким соотношением сторон. Я бы использовал разделяемую память для захвата данных с двумя блоками (каждый idx захватывает 2 элемента из данных, вероятно, разделенных элементами blockDim.x), затем каждый блок выполняет свои назначенные «нечетные» строки, за которыми следуют «четные» строки. Вам придется перекомпоновать ix и iy (и, вероятно, idx), и вы будете использовать 1/2 как можно больше блоков, но должен быть объединенный доступ к памяти, за которым следует недивергентный код.

Answer 3

Разнообразие потоков здесь не является большой проблемой по сравнению с доступом с чередующейся памятью с точки зрения производительности. Я собирался объединиться. Кроме того, ваше хранилище данных имеет неявное упорядочение AoS. Если вы можете изменить порядок данных в SoA, вы можете решить обе проблемы.

Таким образом, я бы изменил порядок этого ядра, чтобы сначала обрабатывать вещи по-крупному. Это решает проблему коалесценции, но вносит разницу в деформации.

Если вы не можете переупорядочить данные, я бы подумал об устранении детерминант warp, изменив схему индексирования, так что даже warps обрабатывают четные элементы, а нечетные искажения обрабатывают нечетные элементы.

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