2011-01-13 106 views
2

我有两个计算类似内容的CUDA内核。一种是使用全局内存(myfun是一种从全局内存中读取很多内容并进行计算的设备函数)。第二个内核将该块数据从全局内存传输到共享内存,以便数据可以在块的不同线程之间共享。使用全局内存的内核比共享内存的内核快得多。可能的原因是什么?CUDA中的全局vs共享内存

loadArray只是将d_x的一小部分复制到m

__global__ void mykernel(float *d_x, float *d_y, int *d_z, float *d_u, int N, int K, int D) 
{ 

    int tid = blockIdx.x*blockDim.x + threadIdx.x; 
    int index = 0; 
    float max_s = 1e+37F; 


    if (tid < N) 
    { 

     for (int i = 0; i < K; i++) 
     { 

      float s = myfun(&d_x[i*D], d_y, tid); 

      if (s > max_s) 
      { 
       max_s = s; 
       index = i; 
      } 
     } 

     d_z[tid] = index; 
     d_u[tid] = max_s; 
    } 
} 

使用共享内存:

__global__ void mykernel(float *d_x, float *d_y, int *d_z, float *d_u, int N, int K) 
{ 
    int tid = blockIdx.x*blockDim.x + threadIdx.x; 
    int index = 0; 
    float max_s = 1e+37F; 

    extern __shared__ float m[]; 
    if(threadIdx.x == 0) 
    loadArray(m, d_x); 
    __syncthreads(); 

    if (tid < N) 
    { 

     for (int i = 0; i < K; i++) 
     { 

      float s = myfun(m, d_y, tid); 

      if (s > max_s) 
      { 
       max_s = s; 
       index = i; 
      } 
     } 

     d_z[tid] = index; 
     d_u[tid] = max_s; 
    } 
} 

回答

3

的问题是,只有在每一个块中的第一个线程从全局内存读取到共享内存中,这是不是让从全局内存读取的所有线程慢得多同时。

当单个线程需要访问全局内存中的相邻元素时,使用共享内存是一个优势 - 但这似乎并不是这种情况。

0

IMO,如果您有平行nsight安装在说Windows机器和处决进行跟踪,你可能有更多的见解。或者,通过您的应用运行cudaprof以尝试找出可能的延迟时间。