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