作者lgen7604 ()
看板C_and_CPP
標題Re: [問題] 使用CUDA Share memory執行"點乘"比glo …
時間Thu Mar 18 02:38:36 2010
主要是你寫的 Matrix_Point_Multiplication_SM 有問題
你用for loop的功能是為了什麼?
能夠只執行一次的動作應該不需要重複執行好幾次吧
我把 Matrix_Point_Multiplication_SM 改成如下
呼叫的時候用
Matrix_Point_Multiplication_SM<<<blocks, threads,
BLOCK_SIZE*BLOCK_SIZE*2*sizeof(float)>>>( d_input1, d_input2, d_output_s );
// ------------------- 用 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 ty = threadIdx.y;
int tx = threadIdx.x;
int row = blockIdx.y * BLOCK_SIZE ;
int col = blockIdx.x * BLOCK_SIZE ;
if( (row+ty < NNy) && (col+tx < NNx) )
{
// 矩陣搬移Share memory
Msm[ ty ][ tx ] = Ma[ (row+ty)*NNx + (col+tx) ];
Nsm[ ty ][ tx ] = Nb[ (row+ty)*NNx + (col+tx) ];
// 矩陣運算Share memory
Pc[ (row+ty)*NNx + (col+tx) ] = Msm[ ty ][ tx ] * Nsm[ ty ][ tx ];
}
}
--
--
※ 發信站: 批踢踢實業坊(ptt.cc)
◆ From: 122.120.44.12
推 aada:我是從矩陣乘法改過來的, 少考慮了會重複的問題 03/18 03:02
推 aada:為什麼block_size*block_size之後要乘2呢 03/18 03:09
→ lgen7604:因為有Msm和Nsm啊 03/18 03:12
推 hardcover:和在excution configuration指定s.m.大小應該無關 03/18 03:14
→ lgen7604:我記得使用到s.m.但呼叫kernel時沒有加入size參數的話 03/18 03:18
→ lgen7604:可能會有無法預期的錯誤 03/18 03:19
推 hardcover:可以說一下是哪裡說的嗎? programming guide 有講指定 03/18 03:21
→ hardcover:第三個參數的時機吧? 好像是你沒定大小的時候? 03/18 03:22
→ lgen7604:學長之前做實驗碰到的Bug 我翻過programming guide是沒有 03/18 03:24
→ lgen7604:提到使用s.m.但是沒加參數會怎麼樣 03/18 03:25
推 hardcover:如果有指定大小，正常應該是不是加才對 03/18 03:30
→ lgen7604:找好久 終於找到了 Appendix.B.12 沒有寫的話預設的 03/18 03:35
→ lgen7604:shared memory size是0 03/18 03:35
推 hardcover:嗯，是dynamic allocate時用的，所以是runtime指定大小 03/18 03:38
→ hardcover:原來的code已經在compile time指定大小了，所有有兩份? 03/18 03:39
→ hardcover:好像也沒有，因為沒有宣告? (extern __shared__ xxx) 03/18 03:40