Блог

Я оценил две производительности ядра:

 #include <chrono>
#include <cuda_runtime.h>
#include <stdio.h>

void initData_int(int *p, int size){
    for (int t=0; t<size; t  ){
        p[t] = (int)(rand()amp;0xff);
    }
}

__global__ void reduceShfl(int *in, int* out, int size)
{
    extern __shared__ int smem[];
    int tid = threadIdx.x;
    int idx = threadIdx.x   blockIdx.x*blockDim.x*4;
    smem[tid] = 0;
    if (tid>=size) return;
    int tmp = 0; 
    if (idx   blockDim.x*3 <= size){
        int a = in[idx];
        int b = in[idx blockDim.x];
        int c = in[idx 2*blockDim.x];
        int d = in[idx 3*blockDim.x];
        tmp = a   b   c   d;
    }
    smem[tid] = tmp;
    __syncthreads();

    if (blockDim.x >= 1024 amp;amp; tid < 512){
        smem[tid]  = smem[tid   512];
    }
    __syncthreads();
    if (blockDim.x >= 512 amp;amp; tid < 256){
        smem[tid]  = smem[tid   256];
    }
    __syncthreads();
    if (blockDim.x >= 256 amp;amp; tid < 128){
        smem[tid]  = smem[tid   128];
    }
    __syncthreads();
    if (blockDim.x >= 128 amp;amp; tid < 64){
        smem[tid]  = smem[tid   64];
    }
    __syncthreads();

    if (blockDim.x >= 64 amp;amp; tid < 32){
        smem[tid]  = smem[tid   32];
    }
    __syncthreads();
    
    int tmpsum = smem[tid]; 
    tmpsum  = __shfl_xor_sync(0xffffffff, tmpsum, 16);
    tmpsum  = __shfl_xor_sync(0xffffffff, tmpsum, 8);
    tmpsum  = __shfl_xor_sync(0xffffffff, tmpsum, 4);
    tmpsum  = __shfl_xor_sync(0xffffffff, tmpsum, 2);
    tmpsum  = __shfl_xor_sync(0xffffffff, tmpsum, 1);
    if (tid==0)
        out[blockIdx.x] = tmpsum;
}

__global__ void reduceShmUnroll(int *in, int *out, int num)
{
    extern __shared__ int smem[];
    int tid = threadIdx.x;
    int idx = threadIdx.x   blockIdx.x*blockDim.x*4;
    
    if (tid >= num) return;
    int tmp=0;
    if(idx   blockDim.x*3 <= num)
    {
        int a = in[idx];
        int b = in[idx   blockDim.x];
        int c = in[idx   blockDim.x*2];
        int d = in[idx   blockDim.x*3];
        tmp = a   b   c   d;
    }   
    smem[tid] = tmp;
    __syncthreads();

    if (blockDim.x >= 1024 amp;amp; tid < 512){
        smem[tid]  = smem[tid   512];
    }
    __syncthreads();
    if (blockDim.x >= 512 amp;amp; tid < 256){
        smem[tid]  = smem[tid 256];
    }
    __syncthreads();
    if (blockDim.x >= 256 amp;amp; tid < 128){
        smem[tid]  = smem[tid 128];
    }
    __syncthreads();
    if (blockDim.x >= 128 amp;amp; tid < 64){
        smem[tid]  = smem[tid 64];
    }
    __syncthreads();
    if (tid < 32){
        volatile int *vsmem = smem;
        vsmem[tid]  = vsmem[tid 32];
        vsmem[tid]  = vsmem[tid 16];
        vsmem[tid]  = vsmem[tid 8];
        vsmem[tid]  = vsmem[tid 4];
        vsmem[tid]  = vsmem[tid 2];
        vsmem[tid]  = vsmem[tid 1];
    }

    if (tid == 0) out[blockIdx.x] = smem[0];
}

int main(int agrc, char **argv)
{
    int size = 1<<24;
    int nBytes = size*sizeof(int);
    int *a_h = (int*)malloc(nBytes);
    initData_int(a_h, size);

    int blocksize = 1024;
    int gridsize = (size-1)/blocksize 1;
    dim3 block(blocksize, 1);
    dim3 grid((size-1)/blocksize 1, 1);
    int *a_d, *b_d;
    cudaMalloc((int**)amp;a_d, nBytes);
    cudaMalloc((int**)amp;b_d, grid.x*sizeof(int));
    cudaMemcpy(a_d, a_h, nBytes, cudaMemcpyHostToDevice);
    int *tmp = (int*)malloc(gridsize*sizeof(int));
    memset(tmp, 0, grid.x/4);
    cudaMemcpy(a_d, a_h, nBytes, cudaMemcpyHostToDevice);
    auto s_0 = std::chrono::system_clock::now();
    reduceShfl<<<grid, block, blocksize*sizeof(int)>>>(a_d, b_d, size);
    cudaMemcpy(tmp, b_d, grid.x/4*sizeof(int), cudaMemcpyDeviceToHost);
    cudaDeviceSynchronize();
    int res_1 = 0;
    for (int i=0; i<grid.x/4; i  ){
        res_1  = tmp[i];
    }
    auto e_0 = std::chrono::system_clock::now();
    std::chrono::duration<double> diff = e_0 - s_0;
    printf("Result from reduceShfl is: %d and time cost is /.n", res_1, diff.count());

    memset(tmp, 0, grid.x/4);
    cudaMemcpy(a_d, a_h, nBytes, cudaMemcpyHostToDevice);
    s_0 = std::chrono::system_clock::now();
    reduceShmUnroll<<<grid, block, blocksize*sizeof(int)>>>(a_d, b_d, size);
    cudaMemcpy(tmp, b_d, grid.x/4*sizeof(int), cudaMemcpyDeviceToHost);
    cudaDeviceSynchronize();
    int res_0 = 0;
    for (int i=0; i<grid.x/4; i  ){
        res_0  = tmp[i];
    }
    e_0 = std::chrono::system_clock::now();
    diff = e_0 - s_0;
    printf("Result from reduceShmUnroll is: %d and time cost is /.n", res_0, diff.count());

    cudaFree(a_d);
    cudaFree(b_d);
    free(a_h);
    free(tmp);
    return 0;
}
 

Основное различие заключается в последнем уменьшении деформации, reduceShmUnroll использовании общей памяти и reduceShfl выполнении перетасовки деформации, которая использует регистровую память.
Но я обнаружил, что reduceShfl это медленнее, чем reduceShmUnroll .

 Result from reduceShfl is: 2139353471 and time cost is 0.000533.
Result from reduceShmUnroll is: 2139353471 and time cost is 0.000485.
 

Что-то не так с моим кодом?

Что-то не так с моим кодом?

Да, я бы сказал, что с вашим кодом что-то не так.

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

общее ядро mem:

 __syncthreads();
if (tid < 32){  // this is missing from your shuffle kernel
    volatile int *vsmem = smem;
    vsmem[tid]  = vsmem[tid 32];
    vsmem[tid]  = vsmem[tid 16];
    vsmem[tid]  = vsmem[tid 8];
    vsmem[tid]  = vsmem[tid 4];
    vsmem[tid]  = vsmem[tid 2];
    vsmem[tid]  = vsmem[tid 1];
}
 

перетасовать ядро:

 __syncthreads();

int tmpsum = smem[tid]; 
tmpsum  = __shfl_xor_sync(0xffffffff, tmpsum, 16);
tmpsum  = __shfl_xor_sync(0xffffffff, tmpsum, 8);
tmpsum  = __shfl_xor_sync(0xffffffff, tmpsum, 4);
tmpsum  = __shfl_xor_sync(0xffffffff, tmpsum, 2);
tmpsum  = __shfl_xor_sync(0xffffffff, tmpsum, 1);
if (tid==0)
    out[blockIdx.x] = tmpsum;
 

Когда я ограничиваю ваше ядро shuffle так же, как ограничено общее ядро mem (чтобы ненужные искажения не выполняли ненужную работу), я наблюдаю примерно одинаковое время выполнения (разница примерно в 1%) между двумя ядрами, когда я профилирую с nvprof на V100:

                 0.38%  222.76us         1  222.76us  222.76us  222.76us  reduceShmUnroll(int*, int*, int)
                0.37%  220.55us         1  220.55us  220.55us  220.55us  reduceShfl(int*, int*, int)
 

Именно этого я и ожидал. Для такого ограниченного использования нет оснований полагать, что совместное использование mem или перетасовка будут быстрее или медленнее.

Как активность в общей памяти, так и активность в случайном перемещении имеют ограничения по пропускной способности. Поэтому трудно предсказать, что будет быстрее, потому что это зависит от того, что еще происходит в вашем коде. Если ваш код привязан к пропускной способности общей памяти, и вы преобразуете часть этой активности в перетасовку с деформацией, вы, скорее всего, увидите выгоду от перетасовки с деформацией. То же самое утверждение можно сделать и в другом направлении. Для этого конкретного раздела этого конкретного кода, когда он написан правильно/сопоставимо/эквивалентно, вы не связаныпо-разному по пропускной способности общей памяти или по пропускной способности в случайном порядке, поэтому правильное ожидание не имеет разницы в производительности, заменяя одно другим.