靜態與動態CUDA共享內存分配的性能

Question

我有2內核，它們完全相同。其中一個靜態分配共享內存，另一個在運行時動態分配內存。我使用共享內存作爲二維數組。所以對於動態分配，我有一個計算內存位置的宏。現在，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); 

} 

可能的原因是由於這樣的事實兩個內核的反彙編代碼是完全相同的，即使將替換爲int N = rand();也不會改變。