[問題] CUDA Memcpy相關問題

看板C_and_CPP作者 (阿哩他命EX PLUS)時間7年前 (2017/07/12 17:25), 7年前編輯推噓1(1推 0噓 25→)
留言26則, 3人參與, 最新討論串1/1
開發平台(Platform): (Ex: Win10, Linux, ...) Linux GPGPU-Sim 問題(Question)： 目前正在測試dependent kernel kernel A 負責兩個陣列相加存到result array kernel B 把kernel A的reult + 1 再將結果與CPU運算結果比對 一般的寫法是： 1.Memcpy兩個陣列的data到device 2.launch kernel A 3.將result複製回host 4.把result傳到device 5.launch kernel B 6.將運算結果複製回host 7.比對 但我把第三四步驟省略 比對結果也是正確的 請問是為什麼呢？ 是因為kernel B需要的data與kernel A的result array記憶體位置相同？ 或是有其他原因？ 程式碼如下 https://gist.github.com/anonymous/5d8b6c58ce7ecd0407f6595d41fd8a2c 後來我改變他的data_size 例如200(原本是100) 結果前100個正確 後100個錯誤 這又是為什麼？ -- ※ 發信站: 批踢踢實業坊(ptt.cc), 來自: 140.118.155.192 ※ 文章網址: https://www.ptt.cc/bbs/C_and_CPP/M.1499851515.A.90A.html
→
07/12 17:41, , 1F
第三和第四步驟 (36,37 兩行) 本來就是多餘的
07/12 17:41, 1F
→
07/12 17:44, , 2F
3和4感覺沒有必要，可以直接把gpu的pointer傳給下個ker
07/12 17:44, 2F
→
07/12 17:44, , 3F
nel
07/12 17:44, 3F
→
07/12 17:45, , 4F
錯誤的原因... 你問我我問誰...
07/12 17:45, 4F
→
07/12 17:45, , 5F
VectorAdd2 到底實際作用是什麼只有你知道
07/12 17:45, 5F
→
07/12 17:46, , 6F
但我覺得 17 行那個 idx 值可能跟你想的不一樣
07/12 17:46, 6F
→
07/12 17:46, , 7F
試試 int idx = blockIdx.x*blockDim.x + threadIdx.x;
07/12 17:46, 7F
為什麼3,4是多餘的？ 是遇到要使用相同的data就可以不用這兩個步驟的意思？ VectorAdd2作用是把VectorAdd1的結果再加一 整個程式目的是要測試不做3,4會不會影響結果 改成blockIdx.x*blockDim.x + threadIdx.x確實是正確了 不過我block只有一個 那麼blockIdx.x為0 跟原本只用threadIdx.x有差別嗎？ 還是我哪裡誤會了？
→
07/12 18:46, , 8F
你的 block size 是不是高於 100 卻不到 200
07/12 18:46, 8F
你是指VectorAdd2<<<1, size>>>(....) ^^^這個嗎 我block size設跟data size相同 原本100 結果正確 換成200後錯一半
→
07/12 19:05, , 9F
這麼說好了，你把 blockDim.x 存進 result array
07/12 19:05, 9F
→
07/12 19:05, , 10F
再 printf 出來看看和你想的一樣不一樣...
07/12 19:05, 10F
blockIdx.x 印出來是0 blockDim.x 印出來是200 也就是data_size threadIdx.x 印出來是0~199 跟我想的一樣.... 現在有點搞混了 我把idx那行改回我最原本那樣 結果是正確的 我再看清楚我到底改了哪裡 -------------------------------------- 我比較想了解的地方是 有相依的kernel間的data 可以不用回傳? 請問原理是什麼?
→
07/12 19:57, , 11F
這個問題嘛，只是你腦筋打結而已...
07/12 19:57, 11F
→
07/12 19:57, , 12F
device memory 又還沒有 free 掉，本來就不會平白消失
07/12 19:57, 12F
→
07/12 19:58, , 13F
跟前後 kernel 有沒有相依一點關係也沒有
07/12 19:58, 13F
→
07/12 20:01, , 14F
是說你害我很想買新顯示卡.....
07/12 20:01, 14F
0.0 哈哈 犒賞自己 買下去! 不過回到我問的 記憶體沒有free掉 所以意思是他存取同個地方囉 那麼 會存取同位址 是為什麼? 這個可以解釋一下嗎
→
07/12 21:20, , 15F
位址不就存在 dev_A, dev_B ... 裡面，位址值有變嗎?
07/12 21:20, 15F
→
07/12 21:20, , 16F
沒變的話為什麼不是存取同位址
07/12 21:20, 16F
你的意思是 因為都是dev_C 記憶體位置就是那裏 所以我launch VectorAdd2 後面帶參數dev_C 那麼他就會去讀這個記憶體位址 所以就會拿到正確的資料 是這樣嗎?
→
07/12 21:24, , 17F
而且 device 和 host 之間的 memcpy 非常之慢
07/12 21:24, 17F
→
07/12 21:27, , 18F
資料可以直接留給下一個 kernel 何必再傳回來傳過去
07/12 21:27, 18F
很慢我知道 只是我不知道原來可以這樣做
→
07/13 08:45, , 19F
是的，從 allocate 到 free 之間那塊記憶體都固定給你用
07/13 08:45, 19F
→
07/13 08:45, , 20F
所以你負責記住那個位址(放在dev_C)並且用完要負責free
07/13 08:45, 20F
→
07/13 08:46, , 21F
跟平常寫 C 語言使用 host memory 的狀況是一樣的
07/13 08:46, 21F
→
07/13 08:48, , 22F
kernel 憑著 dev_C 的位址就可以存取到同一塊記憶體
07/13 08:48, 22F
好的 這部分我瞭解了 感謝你 ※ 編輯: v00623 (140.118.155.204), 07/13/2017 10:16:46
推
07/16 14:04, , 23F
加個cudaThreadSynchronize();確保GPU和CPU之間的
07/16 14:04, 23F
→
07/16 14:05, , 24F
data consistency。你變成data size變成200，有可能是GPU
07/16 14:05, 24F
→
07/16 14:06, , 25F
算太久，你又沒保證data consistency，所以CPU讀到不正確
07/16 14:06, 25F
→
07/16 14:06, , 26F
的資料
07/16 14:06, 26F
‣ 更多 v00623 的文章
文章代碼(AID): #1PPUhxaA (C_and_CPP)