使用CUDA Share memory執行"點乘"比global memory慢

[taker1986]

請問一下,

我想使用Share memory來執行點乘,

但我測出來, 竟然是使用global memory比較快,

請問為什麼呢?是我的share memory寫錯了嗎

是因為矩陣點乘的重覆率不高嗎(一點對一點乘,並沒有重複)

以下是我的kernel function


// ------------------- 用 share memory ------------------- //
__global__ void Matrix_Point_Multiplication_SM( float *Ma, float *Nb, float *Pc )
{
__shared__ float Msm[ BLOCK_SIZE ][ BLOCK_SIZE ];
__shared__ float Nsm[ BLOCK_SIZE ][ BLOCK_SIZE ];
int k, w;
int ty = threadIdx.y; //0*16,1*16,~15*16
int tx = threadIdx.x; // 0~15 *15
int row = blockIdx.y * BLOCK_SIZE ; //0*512 16*512 32*512...
int col = blockIdx.x * BLOCK_SIZE ; //(0:16:512) *512

for( k=0; k<NNx; k+=BLOCK_SIZE )
{
// 矩陣搬移Share memory
Msm[ tx ][ ty ] = Ma[ (row+tx)*NNx +(k+ty) ];
Nsm[ tx ][ ty ] = Nb[ (row+tx)*NNx +(k+ty) ];
__syncthreads();

// 矩陣運算Share memory
Pc[ (row+tx)*NNx +(k+ty) ] = Msm[ tx ][ ty ] * Nsm[ tx ][ ty ];
__syncthreads();
}
}

// ------------------- 用 global memory ------------------- //
__global__ void Matrix_Point_Multiplication_GM( float *input1, float *input2, float *output1 )
{
int col = blockIdx.x*blockDim.x+threadIdx.x; // x 表示 column
int row = blockIdx.y*blockDim.y+threadIdx.y; // y 表示 row
if( col < NNx && row < NNy )
{
output1[ row*NNx+col ] = input1[ row*NNx+col ] * input2[ row*NNx+col ] ; // 取右邊矩陣
}
}

[无心人_过过小日子]

1)矩陣搬移Share memory中把tx和ty反了,导致gmem访问没结合.
2)
__shared__ float Msm[ BLOCK_SIZE ][ BLOCK_SIZE ];
__shared__ float Nsm[ BLOCK_SIZE ][ BLOCK_SIZE ];
中,第二维为BLOCK_SIZE有bank conflict,改为BLOCK_SIZE+1(最后一个不用)就会好很多.
3)
矩陣搬移Share memory只要执行一次就行.不用放在循环中.

[cuda2010]

"矩陣搬移Share memory只要执行一次"恐怕不行吧？
刚才又看了一遍楼主的程序，这个是矩阵“点乘”，应该属带宽瓶颈型程序（每个元素只用到一次）。这样的话shared mem用处不大，估计即使没有bank conflict速度也比不上后者。

[cuda2010]

我估计bank conflict倒还不是主要原因，第一段代码里面似乎把tx和ty弄反了，造成mem read效率低下。

[OpenHero]

shared memory的时候，有bank conflict，还有过多的数据传递；
而global memory是一次访问，一次读写，而且都是对齐的，当然效率会高一些；

[wildcat1020]

jiadianfen

[wangpichao]

mapped memory 只适合小数据量的，大数据量的不占优势。