2011-01-13 148 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以嘗試找出可能的延遲時間。