Re: [問題] 使用CUDA來擷取矩陣中部分值

看板C_and_CPP (C/C++)作者時間16年前 (2010/02/24 02:01), 編輯推噓0(0015)
留言15則, 2人參與, 最新討論串7/10 (看更多)

如果我今天矩陣是比較大的畫,如512*512

dim3 blocks, threadsIM5好像就不可以這樣寫了,是嗎
沒錯

int bx = (NNx + BLOCK_SIZE - 1) / BLOCK_SIZE;

dim3 blocks(bx, bx);

dim3 threads(BLOCK_SIZE, BLOCK_SIZE);

BLOCK_SIZE 我設為16

我IM5大致上試出來了

dim3 blocksIM5( 16, 73);

dim3 threadsIM5( 32, 7 );
說實話我完全不知道你的blocks和threads為什麼要取這種數字啊 要是矩陣的大小改了怎麼辦?

movaIM5<<<blocksIM5, threadsIM5>>>( d_input, d_IM5 );

這樣剛好處理(16*32=512, 73*7=511) --> 512*511的矩陣

我這樣寫的方式OK嗎,有什麼地方可以改進的
從你乘出來的結果看也沒有好像處理完512*512的矩陣啊 下面是我把你之前的範例改成任意大小的矩陣都能處理的type NNx NNy BLOCK_SIZE 都可以任意設 (只要沒有overflow) #include <stdio.h> #include <stdlib.h> #include <string.h> #include <cuda_runtime.h> #include <cutil.h> #define NNx 16 #define NNy 16 #define BLOCK_SIZE 8 #define ALIGN 3 // 對齊printf用 __global__ void movaIM5(float *input1, float *output1) { int col = blockIdx.x*blockDim.x+threadIdx.x; int row = blockIdx.y*blockDim.y+threadIdx.y; if(col < NNx && row+1 < NNy) output1[row*NNx+col] = input1[(row+1)*NNx+col]; // 取下面矩陣 //if(col < NNx && row < (NNy-1)) // output1[row*NNx+col] = input1[row*NNx+col]; // 取上面矩陣 } __global__ void movaIM6(float *input1, float *output1) { int col = blockIdx.x*blockDim.x+threadIdx.x; int row = blockIdx.y*blockDim.y+threadIdx.y; if(col+1 < NNx && row < NNy) output1[row*(NNx-1)+col] = input1[row*NNx+col+1]; // 取右邊矩陣 //if(col < (NNx-1) && row < NNy) // output1[row*(NNx-1)+col] = input1[row*NNx+col]; // 取左邊矩陣 } int main(int argc, char* argv[]) { int i; float input[NNx*NNy]; float *d_input; for(i = 0; i < NNx*NNy; i++){ input[i] = i+1; if(i % NNx == 0) printf("\n"); else printf(" "); printf("%*.0f", ALIGN, input[i]); } printf("\n"); cudaMalloc((void**)&d_input, sizeof(float)*NNx*NNy); cudaMemcpy( d_input, input, sizeof(float)*NNx*NNy, cudaMemcpyHostToDevice ); float *d_IM5, *d_IM6; cudaMalloc((void**)&d_IM5, sizeof(float)*NNx*(NNy-1)); cudaMalloc((void**)&d_IM6, sizeof(float)*(NNx-1)*NNy); int bx = (NNx + BLOCK_SIZE - 1) / BLOCK_SIZE; int by = (NNy + BLOCK_SIZE - 1) / BLOCK_SIZE; dim3 blocks(bx,by); dim3 threads(BLOCK_SIZE,BLOCK_SIZE); movaIM5<<<blocks, threads>>>( d_input, d_IM5 ); movaIM6<<<blocks, threads>>>( d_input, d_IM6 ); float *IM5, *IM6; IM5 = (float*) malloc( sizeof(float)*NNx*(NNy-1) ); IM6 = (float*) malloc( sizeof(float)*(NNx-1)*NNy ); cudaMemcpy( IM5, d_IM5, sizeof(float)*NNx*(NNy-1), cudaMemcpyDeviceToHost ); cudaMemcpy( IM6, d_IM6, sizeof(float)*(NNx-1)*NNy, cudaMemcpyDeviceToHost ); for(i = 0; i < NNx*(NNy-1); i++){ if(i % NNx == 0) printf("\n"); else printf(" "); printf("%*.0f", ALIGN, IM5[i]); } printf("\n"); for(i = 0; i < (NNx-1)*NNy; i++){ if(i % (NNx-1) == 0) printf("\n"); else printf(" "); printf("%*.0f", ALIGN, IM6[i]); } printf("\n"); cudaFree(d_input); cudaFree(d_IM5); cudaFree(d_IM6); system("pause"); return 0; } -- ※ 發信站: 批踢踢實業坊(ptt.cc) ◆ From: 122.120.40.234

02/24 02:26, , 1F
我那樣設定是為了我只是取512*511的矩陣,所以將block,threds
02/24 02:26, 1F

02/24 02:27, , 2F
設為那樣, 如果矩陣不一樣的話,就得改了
02/24 02:27, 2F

02/24 07:10, , 3F
input矩陣是512*512還算親切的大小 (511可以分解為73*7)
02/24 07:10, 3F

02/24 07:13, , 4F
如果是510*510該怎麼處理才好? 像這種case就是做參數化時
02/24 07:13, 4F

02/24 07:14, , 5F
要特別注意的地方了
02/24 07:14, 5F
※ 編輯: lgen7604 來自: 122.120.40.234 (02/24 07:22)

02/24 14:17, , 6F
謝謝你, 想請教你一個觀念問題, 你怎麼想出來是用if來寫呢,
02/24 14:17, 6F

02/24 14:19, , 7F
以及寫CUDA時您是怎麼來思考的, 我對CUDA的經驗還沒有很深.
02/24 14:19, 7F

02/25 01:35, , 8F
寫code的想法這個問題 我只能說...經驗吧 看得多想得多
02/25 01:35, 8F

02/25 01:35, , 9F
碰到問題的時候想法自然就會出來了
02/25 01:35, 9F

02/25 01:36, , 10F
寫CUDA要特別注意的就是parallel的觀念了 和一般寫C或C++
02/25 01:36, 10F

02/25 01:36, , 11F
不一樣 要考慮如何把問題分割給很多的block和thread處理
02/25 01:36, 11F

02/25 01:36, , 12F
還有每個block和thread該負責什麼部份 特別是參數化時 要
02/25 01:36, 12F

02/25 01:37, , 13F
注意矩陣的dimension 還有boundary的問題
02/25 01:37, 13F

02/25 01:48, , 14F
如果要考慮performance那問題就多了= = 建議可以看看
02/25 01:48, 14F

02/25 01:49, , 15F
nVIDIA官方的CUDA programming guide 應該會有幫助的
02/25 01:49, 15F
文章代碼(AID): #1BX1V_Sa (C_and_CPP)
討論串 (同標題文章)
文章代碼(AID): #1BX1V_Sa (C_and_CPP)