2014-04-03 234 views
0

我有一個大小爲NxN的複雜數據的矩陣u,我想每行乘以一個大小爲1xN的向量k。 u中的數據按行存儲。cuda中的共享內存

我有兩個這樣的實現。一個利用共享內存,將矩陣分成瓦片,另一個沒有。

我發現共享內存實現multiply1不會更快,並且系統性地一樣快,甚至比multiply2慢。

共享存儲器實現如下,

__global__ void multiply1(cufftComplex *u, cufftComplex *k) { 
    __shared__ cufftComplex k_s[BLOCK_WIDTH]; 
    int idx = blockDim.x*blockIdx.x + threadIdx.x; 
    int idy = blockDim.y*blockIdx.y + threadIdx.y; 
    int index; 

    if (threadIdx.y == 0 && idx < N) { 
     k_s[threadIdx.x] = k[idx]; 
    } 
    __syncthreads(); 

    if (idx < N && idy < N) { 
     index = N*idy + idx; 
     u[index] = cuCmulf(k_s[threadIdx.x],u[index]); 
    } 

} 

鑑於全球存儲器實現如下,

__global__ void multiply2(cufftComplex *u, cufftComplex *k) { 
     int idx = blockDim.x * blockIdx.x + threadIdx.x; 

     if (idx < N*N) { 
      u[idx] =cuCmulf(k[idx % N],u[idx]); 
     } 
    } 

和主函數調用,對於大小的矩陣64×64

dim3 block(16,16); 
dim3 grid(4,4); 
multiply1<<<grid, block>>>(d_u, d_k); 
multiply2<<<16, 256>>>(d_u, d_k); 

如何使用探查器找出爲什麼multiply1不是gett速度至少略有增加?哪些指標可以闡明究竟發生了什麼?

分析器告訴我multiply1,我得到152 GB/s的全局內存加載吞吐量,而multiply2我得到81 GB /秒。這是合乎邏輯的,因爲我從全局內存中加載較少。這不應該轉化爲更快的執行嗎?

回答

1

如果你多次使用它會更快,但在這裏你只用了一次。您在變換你的問題:

copy from global memory to shared memory 
read from shared memory 

代替:

read from global memory 

所以是的,它肯定是比以前的算法只使用全局內存要慢。如果你想利用共享內存,你的算法要多次讀取它,否則你不會花費全局內存。

+0

我從全局存儲器複製每塊一次k的適當的元件,而不是每一次元素,因爲if語句的: 如果(threadIdx.y == 0 && IDX user3495341

+0

每塊一次,但同一塊中的線程正在等待複製完成。 –

+0

我該如何檢查? 會有一種方法可以在探查器中執行此操作嗎? – user3495341