Более быстрое копирование массива при использовании меньшего количества потоков в CUDA

Я протестировал два разных подхода к копированию двумерного массива в ядре CUDA.

Первый запускает блоки потоков TILE_DIM x TILE_DIM. Каждый блок копирует тайл массива, назначая по одному потоку на элемент:

__global__ void simple_copy(float *outdata, const float *indata){

int x = blockIdx.x * TILE_DIM + threadIdx.x;
int y = blockIdx.y * TILE_DIM + threadIdx.y;

int width = gridDim.x * TILE_DIM;

outdata[y*width + x] = indata[y*width + x];

}

Второй взят из блога NVIDIA. Оно похоже на предыдущее ядро, но использует потоки TILE_DIM x BLOCK_ROWS на блок. Каждый поток перебирает несколько элементов матрицы:

__global__ void fast_copy(float *outdata, const float *indata)
{
int x = blockIdx.x * TILE_DIM + threadIdx.x;
int y = blockIdx.y * TILE_DIM + threadIdx.y;
int width = gridDim.x * TILE_DIM;

for (int k = 0 ; k < TILE_DIM ; k += BLOCK_ROWS)
    outdata[(y+k)*width + x] = indata[(y+k)*width + x];
}

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

Визуальный профилировщик NVIDIA подтверждает этот тест.

Так как же второму ядру удается добиться более быстрой копии?

Это полный код, который я использовал для тестирования ядер:

#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>
#include <conio.h>

#define TILE_DIM 32
#define BLOCK_ROWS 8

/* KERNELS */

__global__ void simple_copy(float *outdata, const float *indata){

int x = blockIdx.x * TILE_DIM + threadIdx.x;
int y = blockIdx.y * TILE_DIM + threadIdx.y;

int width = gridDim.x * TILE_DIM;

outdata[y*width + x] = indata[y*width + x];

}
//###########################################################################

__global__ void fast_copy(float *outdata, const float *indata)
{
int x = blockIdx.x * TILE_DIM + threadIdx.x;
int y = blockIdx.y * TILE_DIM + threadIdx.y;
int width = gridDim.x * TILE_DIM;

for (int k = 0 ; k < TILE_DIM ; k += BLOCK_ROWS)
    outdata[(y+k)*width + x] = indata[(y+k)*width + x];
}
//###########################################################################

/* MAIN */

int main(){

float *indata,*dev_indata,*outdata1,*dev_outdata1,*outdata2,*dev_outdata2;
cudaEvent_t start, stop;
float time1,time2;
int i,j,k;

int n_iter = 100;

int N = 2048;

cudaEventCreate(&start);
cudaEventCreate(&stop);


dim3 grid(N/TILE_DIM, N/TILE_DIM);
dim3 threads1(TILE_DIM,TILE_DIM);
dim3 threads2(TILE_DIM,BLOCK_ROWS);

// Allocations

indata = (float *)malloc(N*N*sizeof(float));
outdata1 = (float *)malloc(N*N*sizeof(float));
outdata2 = (float *)malloc(N*N*sizeof(float));

cudaMalloc( (void**)&dev_indata,N*N*sizeof(float) );
cudaMalloc( (void**)&dev_outdata1,N*N*sizeof(float) );
cudaMalloc( (void**)&dev_outdata2,N*N*sizeof(float) );

// Initialisation

for(j=0 ; j<N ; j++){
        for(i=0 ; i<N ; i++){
            indata[i + N*j] = i + N*j;
        }
}

// Transfer to Device
cudaMemcpy( dev_indata, indata, N*N*sizeof(float),cudaMemcpyHostToDevice );

// Simple copy
cudaEventRecord( start, 0 );
for(k=0 ; k<n_iter ; k++){
    simple_copy<<<grid, threads1>>>(dev_outdata1,dev_indata);
}
cudaEventRecord( stop, 0 );

cudaEventSynchronize( stop );
cudaEventElapsedTime( &time1, start, stop );
printf("Elapsed time with simple copy: %f\n",time1);

// Fast copy
cudaEventRecord( start, 0 );
for(k=0 ; k<n_iter ; k++){
    fast_copy<<<grid, threads2>>>(dev_outdata2,dev_indata);
}
cudaEventRecord( stop, 0 );

cudaEventSynchronize( stop );
cudaEventElapsedTime( &time2, start, stop );
printf("Elapsed time with fast copy: %f\n",time2);

// Transfer to Host

cudaMemcpy( outdata1, dev_outdata1, N*N*sizeof(float),cudaMemcpyDeviceToHost );
cudaMemcpy( outdata2, dev_outdata2, N*N*sizeof(float),cudaMemcpyDeviceToHost );

// Check for error
float error = 0;
for(j=0 ; j<N ; j++){
        for(i=0 ; i<N ; i++){
            error += outdata1[i + N*j] - outdata2[i + N*j];
        }
}
printf("error: %f\n",error);

/*// Print the copied matrix
printf("Copy\n");
for(j=0 ; j<N ; j++){
        for(i=0 ; i<N ; i++){
            printf("%f\t",outdata1[i + N*j]);
        }
        printf("\n");
}*/

cudaEventDestroy( start );
cudaEventDestroy( stop );

free(indata);
free(outdata1);
free(outdata2);

cudaFree(dev_indata);
cudaFree(dev_outdata1);
cudaFree(dev_outdata2);

cudaDeviceReset();

getch();

return 0;
 }

//###########################################################################