Re: [心得] 使用CUDA Share memory

[clhsieh ()]

※ 引述《lgen7604 ()》之銘言：
: 根據hardcover大大提供的情報
: 還有之前做實驗遇到Bug的經驗
: 加上 CUDA Programming Guide 2.3 的說明
: (參考Appendix.B.2.3,  B.12, 3.2.2)
: 我覺得 shared memory 的使用應該是這樣
: 根據Appendix.B.2.3及B.12
: 在 dynamic allocate 的時候
: 必須在 excution configuration 指定使用的 shared memory 大小
: 至於已經宣告指定大小的 shared memory
: 在呼叫kernel時可以不必加上指定 shared memory 大小的參數
: 也可以加上指定 shared memory 大小的參數
: 不過必須和宣告的 shared memory 大小一致
: 否則的話可能在程式執行過程中會有問題
: 所以好像不寫比較不容易出錯 = =
: 感謝hardcover大大提供討論與指教
借這個標題問一下
我最近在寫filter
由於想讓filter的size可以動態決定
所以使用share memory的時候
必須allocate [blockDim.x+2*r][block.y+2*r] 來放資料
其中r為filter的size

照上面這樣說的話我就必須在呼叫kernel的時候使用
kernel<<<grid, block, (block.x+2*r)*(block.y+2*r)>>>()
來allocate share memory的空間
這樣的話我的share memory就會是一維的
在讀取資料的時候需要多一些運算去計算位置
有辦法動態宣告出二維的share memory嗎??

--
※ 發信站: 批踢踢實業坊(ptt.cc) 
◆ From: 140.112.29.95
→ clhsieh:阿 我現在想到了
→ clhsieh:在kernel裡面再用指標去指向這空間 這樣應該可行吧?
→ clhsieh:或著是說有沒有其他比較好的方法?
推 LittleBob:我很好奇 用一維宣告的陣列 會因為需要計算而降低效能嘛
→ clhsieh:我也不是很確定 不過算位置要用乘法
→ clhsieh:弄二維的話等於乘法只在一開始就做一次記錄下來了
→ clhsieh:不過我不知道這樣memory acess會不會反而變更久...
推 hardcover:memory不是本來就1D嗎?再多維最後也要map到1D吧?