Re: [問題] CUDA關於shared memory的使用

看板C_and_CPP作者 (資源回收將軍)時間14年前 (2010/08/04 18:25), 編輯推噓2(208)
留言10則, 4人參與, 最新討論串1/1
※ 引述《a88258850 (資源回收將軍)》之銘言: ( *[1m *[m 為色碼,可以按 Ctrl+V 預覽會顯示的顏色 ) ( 未必需要依照此格式,文章條理清楚即可 ) 遇到的問題: (題意請描述清楚) 我希望shared memory宣告的array, 可以讓我當變數傳進另外宣告的devicefunction... 不知道這種作法的問題在哪邊? 以下是我的程式... __device__ void test(float *t) { .............. } __global__ void kernel( float total) { __shared__ float t[25]; int total_task = blockDim.x * gridDim.x; int task_sn = blockIdx.x * blockDim.x + threadIdx.x; for(int k=task_sn; k <total ;k+=total_task) { test(t); } } int main() { kernel<<<(1,Threads)>>>(total) } 程式跑出來的錯誤結果: 假如我假設Threads = 32,那程式可以正常跑完沒問題。 但,Threads >=64以後,程式會卡住... 開發平台: (例: VC++ or gcc/g++ or Dev-C++, Windows or Linux) linux 板子是Tesla C2050,每個block有49152 bytes的shared memory; 若設定64個thread,每個thread大概有768bytes = 192個float, 謝謝幫忙解答。 -- ※ 發信站: 批踢踢實業坊(ptt.cc) ◆ From: 140.113.90.80

08/04 18:09,
如果用global memory的話,程式不會有問題。
08/04 18:09

08/04 18:10,
會不會是 sm 爆掉了?
08/04 18:10

08/04 18:11,
每個 block 的 shared memory 好像有大小限制
08/04 18:11

08/04 18:12,
每塊板子可能都不同 要查一下喔 @@"
08/04 18:12

08/04 18:17,
實驗室是Tesla,shared memory有49152 bytes/block。
08/04 18:17
-- ※ 發信站: 批踢踢實業坊(ptt.cc) ◆ From: 140.113.90.80

08/04 18:31, , 1F
可以用大E修文 原文看要不要清掉
08/04 18:31, 1F

08/04 21:14, , 2F
google "shared memory pointer cuda"
08/04 21:14, 2F

08/04 21:15, , 3F
好像在哪篇paper有看到說, 這是 compiler 問題
08/04 21:15, 3F

08/04 21:15, , 4F
因為 ptx code 上 global memory 跟 shared memory 指令不
08/04 21:15, 4F

08/04 21:16, , 5F
一樣, 我們做最壞假設, test 是再另外一個檔案編譯
08/04 21:16, 5F

08/04 21:17, , 6F
那麼 test 就不知道到底你傳的是不是 shared memory ptr
08/04 21:17, 6F

08/04 22:23, , 7F
感謝bill & akasan的回答^^
08/04 22:23, 7F

08/04 22:32, , 8F
test() 不是會inline嗎?那應該能分得出float *t在哪?
08/04 22:32, 8F

08/04 22:57, , 9F
__device__ 並沒有 imply inline, 即使你標 inline
08/04 22:57, 9F

08/04 22:57, , 10F
compiler 也可以不鳥你XD
08/04 22:57, 10F
文章代碼(AID): #1CMK0aFp (C_and_CPP)