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

У меня есть 5 массивов больших размеров A(N*5), B(N*5), C(N*5), D(N*5), E(N*2) номер 5 и 2 представляет компоненты этих переменных в разных плоскостях / осях. Вот почему я структурировал массивы таким образом, чтобы я мог визуализировать данные, когда пишу свой код. N ~ 200^3 ~ 8e06 узлов

Например: это то, как выглядит мое ядро ​​в простейшей форме, где я делаю все вычисления в глобальной памяти.

#define N 200*200*200
__global__ void kernel(doube *A, double *B, double *C, 
            double *D, double *E, double *res1, double *res2, 
            double *res3, double *res4 )
    {
       int a, idx=threadIdx.x + blockIdx.x * blockDim.x;
        if(idx>=N) {return;}
        res1[idx]=0.; res2[idx]=0.; 
        res3[idx]=0.; res4[idx]=0.

        for (a=0; a<5; a++)
        {
            res1[idx] += A[idx*5+a]*B[idx*5+a]+C[idx*5+a] ;
            res2[idx] += D[idx*5+a]*C[idx*5+a]+E[idx*2+0] ;
            res3[idx] += E[idx*2+0]*D[idx*5+a]-C[idx*5+a] ;
            res4[idx] += C[idx*5+a]*E[idx*2+1]-D[idx*5+a] ;
        }

    }

Я знаю, что цикл for можно исключить, но я оставил его здесь, так как на него удобно смотреть код. Это работает, но, очевидно, это крайне неэффективно и медленно для карты Tesla K40 даже после удаления цикла for. Арифметика, показанная внутри цикла for, просто для того, чтобы дать представление, фактические вычисления намного длиннее и запутаны с res1,res2..., также входящими в микс.

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

    #define THREADS_PER_BLOCK 256
    __global__ void kernel_shared(doube *A, double *B, double *C, 
               double *D, double *E, double *res1, double *res2, 
               double *res3, double *res4  )
    {
       int a, idx=threadIdx.x + blockIdx.x * blockDim.x;
       int ix = threadIdx.x;
       __shared__ double A_sh[5*THREADS_PER_BLOCK];
       __shared__ double B_sh[5*THREADS_PER_BLOCK];
       __shared__ double C_sh[5*THREADS_PER_BLOCK];
       __shared__ double D_sh[5*THREADS_PER_BLOCK];
       __shared__ double E_sh[2*THREADS_PER_BLOCK];

       //Ofcourse this will not work for all arrays in shared memory; 
        so I am allowed  to put any 2 or 3 variables (As & Bs) of  
         my choice in shared and leave rest in the global memory. 

       for(int a=0; a<5; a++)
     {
        A_sh[ix*5 + a] = A[idx*5 + a] ;
        B_sh[ix*5 + a] = B[idx*5 + a] ;
     }
            __syncthreads();



    if(idx>=N) {return;}
        res1[idx]=0.; res2[idx]=0.; 
        res3[idx]=0.; res4[idx]=0.
    for (a=0; a<5; a++)
    {
        res1[idx] += A_sh[ix*5+a]*B_sh[ix*5+a]+C[idx*5+a];
        res2[idx] += B_sh[ix*5+a]*C[idx*5+a]+E[idx*2+0]  ;
        res3[idx] += E[idx*2+0]*D[idx*5+a]-C[idx*5+a]    ;
        res4[idx] += B_sh[ix*5+a]*E[idx*2+1]-D[idx*5+a]  ;
    }

}

Это немного помогает, но я хотел бы реализовать один из этих подходов к сокращению (без конфликта в банке) для повышения производительности, где я могу поместить все свои переменные в общий (может быть подход с использованием листов), а затем выполнить расчетную часть. Я видел пример сокращения в папке CUDA_Sample, но этот пример работает для суммирования только по одному вектору без какой-либо сложной арифметики, связанной с несколькими массивами из общей памяти. Я был бы признателен за любую помощь или предложение по улучшению моего существующего подхода kernel_shared, чтобы включить подход сокращения.

Источник

1 ответ

1. Что вам нужно, это не общая память

Изучая ваше исходное ядро, мы замечаем, что для каждого значения aвы используете для суммирования не более 12 значений из четырех дельт (вероятно, менее 12, я точно не считал). Все это прекрасно вписывается в ваш регистровый файл - даже для двойных значений: 12 * sizeof(double), плюс 4 * sizeof(double) для промежуточных результатов составляет 32 4-байтовых регистра на поток. Значительно за предел, даже если у вас есть 1024 потоков на блок.

Теперь причины, по которым ваше ядро ​​работает медленно, в основном

2. Неоптимальные схемы доступа к памяти

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

5 * idx
5 * idx + 1
...
5 * idx + 4

имей ручку

5 * blockDim.x * blockIdx.x + threadIdx.x
5 * blockDim.x * blockIdx.x + threadIdx.x + blockDim.x
...
5 * blockDim.x * blockIdx.x + threadIdx.x + 4 * blockDim.x

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

3. Чрезмерное добавление к местам в глобальной памяти

Эта проблема более специфична для вашего случая. Видите ли, вам действительно не нужно менять resN[idx] значение в глобальном после каждого добавления, и вам, конечно, не нужно читать значение, которое есть там, когда вы собираетесь писать. Когда ваше ядро ​​работает, один поток вычисляет новое значение для resN[idx] - так что он может просто сложить вещи в реестре, и написать resN[idx] когда это будет сделано (даже не глядя на его адрес).

Если вы измените свою схему доступа к памяти, как я предлагал в пункте 1., реализация предложения в пункте 2. становится более сложной, поскольку вам нужно будет сложить значения из нескольких дорожек в одной и той же деформации и, возможно, убедиться, что вы не используете Не пересекайте границы деформации с чтениями, относящимися к одному вычислению. Чтобы узнать, как это сделать, я предлагаю вам взглянуть на эту презентацию о сокращениях в случайном порядке.

Источник
Другие вопросы по тегам