静的対動的CUDA共有メモリ割り当てのパフォーマンス

Question

まったく同じことを行う2のカーネルがあります。そのうちの1つは共有メモリを静的に割り当て、もう1つは実行時に動的にメモリを割り当てます。私は2D配列として共有メモリを使用しています。だから動的割り当てのために、私はメモリの場所を計算するマクロを持っています。さて、2カーネルによって生成された結果はまったく同じです。しかし、私は両方のカーネルから得たタイミングの結果は3回離れています！スタティックメモリの割り当てははるかに高速です。私は自分のコードを投稿できないことを申し訳なく思っています。誰かがこれを正当化することはできますか？

Answer 1

私は、静的共有メモリ割り当てが動的共有メモリ割り当てより高速であるという証拠はありません。上記のコメントで示されているように、再生器なしであなたの質問に答えることは不可能です。静的または動的な共有メモリの割り当てと実行時に以下のコードの少なくとも場合には、同じカーネルのタイミングは、正確に同じである。

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

#define BLOCK_SIZE 512 

/********************/ 
/* CUDA ERROR CHECK */ 
/********************/ 
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); } 
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true) 
{ 
    if (code != cudaSuccess) 
    { 
     fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); 
     if (abort) exit(code); 
    } 
} 

/***********************************/ 
/* SHARED MEMORY STATIC ALLOCATION */ 
/***********************************/ 
__global__ void kernel_static_memory_allocation(int *d_inout, int N) 
{ 
    __shared__ int s[BLOCK_SIZE]; 

    const int tid = threadIdx.x; 
    const int i  = blockIdx.x * blockDim.x + threadIdx.x; 

    if (i < N) { 

     s[tid] = d_inout[i]; 
     __syncthreads(); 

     s[tid] = s[tid] * s[tid]; 
     __syncthreads(); 

     d_inout[i] = s[tid]; 
    } 
} 

/************************************/ 
/* SHARED MEMORY DYNAMIC ALLOCATION */ 
/************************************/ 
__global__ void kernel_dynamic_memory_allocation(int *d_inout, int N) 
{ 
    extern __shared__ int s[]; 

    const int tid = threadIdx.x; 
    const int i  = blockIdx.x * blockDim.x + threadIdx.x; 

    if (i < N) { 

     s[tid] = d_inout[i]; 
     __syncthreads(); 

     s[tid] = s[tid] * s[tid]; 
     __syncthreads(); 

     d_inout[i] = s[tid]; 
    } 
} 

/********/ 
/* MAIN */ 
/********/ 
int main(void) 
{ 
    int N = 1000000; 

    int* a = (int*)malloc(N*sizeof(int)); 

    for (int i = 0; i < N; i++) { a[i] = i; } 

    int *d_inout; gpuErrchk(cudaMalloc(&d_inout, N * sizeof(int))); 

    int n_blocks = N/BLOCK_SIZE + (N%BLOCK_SIZE == 0 ? 0:1); 

    gpuErrchk(cudaMemcpy(d_inout, a, N*sizeof(int), cudaMemcpyHostToDevice)); 

    float time; 
    cudaEvent_t start, stop; 
    cudaEventCreate(&start); 
    cudaEventCreate(&stop); 
    cudaEventRecord(start, 0); 
    kernel_static_memory_allocation<<<n_blocks,BLOCK_SIZE>>>(d_inout, N); 
    gpuErrchk(cudaPeekAtLastError()); 
    gpuErrchk(cudaDeviceSynchronize()); 
    cudaEventRecord(stop, 0); 
    cudaEventSynchronize(stop); 
    cudaEventElapsedTime(&time, start, stop); 
    printf("Static allocation - elapsed time: %3.3f ms \n", time); 

    cudaEventRecord(start, 0); 
    kernel_dynamic_memory_allocation<<<n_blocks,BLOCK_SIZE,BLOCK_SIZE*sizeof(int)>>>(d_inout, N); 
    gpuErrchk(cudaPeekAtLastError()); 
    gpuErrchk(cudaDeviceSynchronize()); 
    cudaEventRecord(stop, 0); 
    cudaEventSynchronize(stop); 
    cudaEventElapsedTime(&time, start, stop); 
    printf("Dynamic allocation - elapsed time: %3.3f ms \n", time); 

} 

そのための可能な理由は、という事実によるものです2つのカーネルの逆アセンブルされたコードはまったく同じで、N = 1000000;をint N = rand();に置き換えても変更されません。