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

[lgen7604 ()]

→ aada:如果我今天矩陣是比較大的畫,如512*512
→ aada:dim3 blocks, threadsIM5好像就不可以這樣寫了,是嗎
沒錯

→ aada:int bx = (NNx + BLOCK_SIZE - 1) / BLOCK_SIZE;
→ aada:dim3 blocks(bx, bx);
→ aada:dim3 threads(BLOCK_SIZE, BLOCK_SIZE);
→ aada:BLOCK_SIZE 我設為16
推 aada:我IM5大致上試出來了
→ aada:dim3 blocksIM5( 16, 73);
→ aada:dim3 threadsIM5( 32, 7 );
說實話我完全不知道你的blocks和threads為什麼要取這種數字啊
要是矩陣的大小改了怎麼辦?

→ aada:movaIM5<<<blocksIM5, threadsIM5>>>( d_input, d_IM5 );
→ aada:這樣剛好處理(16*32=512, 73*7=511) --> 512*511的矩陣
→ aada:我這樣寫的方式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
→ aada:我那樣設定是為了我只是取512*511的矩陣,所以將block,threds
→ aada:設為那樣, 如果矩陣不一樣的話,就得改了
→ lgen7604:input矩陣是512*512還算親切的大小 (511可以分解為73*7)
→ lgen7604:如果是510*510該怎麼處理才好? 像這種case就是做參數化時
→ lgen7604:要特別注意的地方了
※ 編輯: lgen7604        來自: 122.120.40.234       (02/24 07:22)
→ aada:謝謝你, 想請教你一個觀念問題, 你怎麼想出來是用if來寫呢,
→ aada:以及寫CUDA時您是怎麼來思考的, 我對CUDA的經驗還沒有很深.
→ lgen7604:寫code的想法這個問題 我只能說...經驗吧  看得多想得多
→ lgen7604:碰到問題的時候想法自然就會出來了
→ lgen7604:寫CUDA要特別注意的就是parallel的觀念了 和一般寫C或C++
→ lgen7604:不一樣 要考慮如何把問題分割給很多的block和thread處理
→ lgen7604:還有每個block和thread該負責什麼部份 特別是參數化時 要
→ lgen7604:注意矩陣的dimension 還有boundary的問題
→ lgen7604:如果要考慮performance那問題就多了= =  建議可以看看
→ lgen7604:nVIDIA官方的CUDA programming guide  應該會有幫助的