高速化の手法として、1Threadごと1Byteづつ処理しているものを、4Byteまとめて処理するものがあります。
比較的、簡単に高速化ができるのでぜひやってみましょう。
処理内容
8192×8192のBYTE型の二次元配列に対し、縦方向に加算していき、8192の一次元配列に各X座標の加算結果を格納する処理です。
１スレッドが1つのX座標を担当します。
コード
1Byteずつ処理
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 |
//カーネル __global__ void kernel_sum_1byte(BYTE* pSrc, UINT32 *pResult, int sizex, int sizey) { int idx = blockDim.x * blockIdx.x + threadIdx.x; //自分のスレッドxのindex if (idx >= sizex) { return; //配列をはみ出る場合は無視 } unsigned int sum = 0; for (int y = 0; y < sizey; y++) { sum += pSrc[y * sizex + idx]; } //加算結果を格納 pResult[idx] = sum; return; } //カーネル呼び出し関数 void Process_SumGPU_1BYTE() { int sizex = 8192; int sizey = 8192; int length = sizex * sizey; size_t size_src = sizeof(BYTE) * length; size_t size_res = sizeof(UINT32) * sizex; //ホストメモリのポインタ BYTE* pHostSrc; //数値 UINT32* pHostResult; //加算結果 //デバイスメモリのポインタ BYTE* pDevSrc; //数値 UINT32* pDevResult; //加算結果 //ホストメモリの確保 cudaMallocHost(&pHostSrc, size_src); cudaMallocHost(&pHostResult, size_res); //デバイスメモリの確保 cudaMalloc(&pDevSrc, size_src); cudaMalloc(&pDevResult, size_res); //ホストメモリに適当な値を設定 for (int n = 0; n < length; n++) { pHostSrc[n] = 1; } //ホスト->デバイスへ数値を転送 cudaMemcpy(pDevSrc, pHostSrc, size_src, cudaMemcpyHostToDevice); //カーネル dim3 block(128, 1, 1); dim3 grid((sizex + 128 - 1) / 128, 1, 1); kernel_sum_1byte << <grid, block >> > (pDevSrc, pDevResult, sizex, sizey); //デバイス->ホストへ結果を転送 cudaMemcpy(pHostResult, pDevResult, size_res, cudaMemcpyDeviceToHost); //デバイスメモリの開放 cudaFree(pDevSrc); cudaFree(pDevResult); //ホストメモリの開放 cudaFreeHost(pHostSrc); cudaFreeHost(pHostResult); } |
4Byteずつ処理
４Byteずつまとめて処理するコードです。
確保する配列のサイズなどは変わっていませんが、カーネルでは1スレッドが4Byte分まとめて処理しています。
そのため、ブロックサイズなどが1/4になっていたりすることに注意しましょう。
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 |
//カーネル __global__ void kernel_sum_4byte(UINT32* pSrc, UINT32 *pResult, int sizex, int sizey) { int idx = blockDim.x * blockIdx.x + threadIdx.x; //自分のスレッドxのindex if (idx >= sizex) { return; //配列をはみ出る場合は無視 } unsigned int sum0 = 0; unsigned int sum1 = 0; unsigned int sum2 = 0; unsigned int sum3 = 0; //4Byteを1Byteに分解して処理 UINT32 src ; for (int y = 0; y < sizey; y++) { src = pSrc[y * sizex + idx]; sum0 += (src & 0x000000FF); //1Byte目 sum1 += (src & 0x0000FF00) >> 8; //2Byte目 sum2 += (src & 0x00FF0000) >> 16; //3Byte目 sum3 += (src & 0xFF000000) >> 24; //4Byte目 } //加算結果を格納 pResult[idx * 4 + 0] = sum0; pResult[idx * 4 + 1] = sum1; pResult[idx * 4 + 2] = sum2; pResult[idx * 4 + 3] = sum3; return; } //カーネル呼び出し関数 void Process_SumGPU_4BYTE() { int sizex = 8192; int sizey = 8192; int length = sizex * sizey; size_t size_src = sizeof(BYTE) * length; size_t size_res = sizeof(UINT32) * sizex; //ホストメモリのポインタ BYTE* pHostSrc; //数値 UINT32* pHostResult; //加算結果 //デバイスメモリのポインタ BYTE* pDevSrc; //数値 UINT32* pDevResult; //加算結果 //ホストメモリの確保 cudaMallocHost(&pHostSrc, size_src); cudaMallocHost(&pHostResult, size_res); //デバイスメモリの確保 cudaMalloc(&pDevSrc, size_src); cudaMalloc(&pDevResult, size_res); //ホストメモリに適当な値を設定 for (int n = 0; n < length; n++) { pHostSrc[n] = 1; } //ホスト->デバイスへ数値を転送 cudaMemcpy(pDevSrc, pHostSrc, size_src, cudaMemcpyHostToDevice); //カーネル dim3 block(128, 1, 1); dim3 grid((sizex / 4 + 128 - 1) / 128, 1, 1); //4Byte分まとめる //UINT32のポインタにキャスト kernel_sum_4byte <<<grid, block >>> ((UINT32*)pDevSrc, pDevResult, sizex / 4, sizey); //デバイス->ホストへ結果を転送 cudaMemcpy(pHostResult, pDevResult, size_res, cudaMemcpyDeviceToHost); //デバイスメモリの開放 cudaFree(pDevSrc); cudaFree(pDevResult); //ホストメモリの開放 cudaFreeHost(pHostSrc); cudaFreeHost(pHostResult); } |
処理時間
処理時間は以下の通り。
| 1Byteずつ処理 | 533μsec |
| 4Byteまとめて処理 | 356μsec |
なんと、処理時間が3割以上早くなりました。
簡単にできる割には、効果が大きいことが判ります。
