※お役立ち情報満載のOshirabe Goto へ立ち寄ってみてください。(ブックマークしておくと便利です。)

スポンサーサイト 

上記の広告は1ヶ月以上更新のないブログに表示されています。
新しい記事を書く事で広告が消せます。

CUDAの計算 

GPUの設定(前回前々回)ができたところで行列計算を実行。

このサイトを大いに参考にしました。ただし、#include は使いませんでした。(エラーが出たため。)
この行列計算プログラムの意味がわかると、CUDAの使い方が(少し)理解できます。


まずは、3行3列の行列の掛け算 (C=A×B)

cmatrix_3.c:

#include <stdio.h>
#include <malloc.h>
#include <stdlib.h>
#include <time.h>

#define N 3 // 正方行列のサイズを指定(N×N)

int main( int argc, char** argv)
{
int col, row, scan;

int* matA;
int* matB;
int* matC;
clock_t start, end;

// 行列A~Cのメモリ確保
matA = (int*)malloc(sizeof(int) * N*N);
matB = (int*)malloc(sizeof(int) * N*N);
matC = (int*)malloc(sizeof(int) * N*N);


// 行列A、Bの各要素に乱数を入れる
// 行列Cは初期値0
for (col = 0; col < N; col++){
for (row = 0; row < N; row++){
matA[col * N + row] = rand() % 100;
matB[col * N + row] = rand() % 100;
matC[col * N + row] = 0.0;
}
}

start = clock();

// 行列の計算C = A × B
for (col = 0; col < N; col++) {
for (row = 0; row < N; row++) {
for (scan = 0; scan < N; scan++) {
matC[col * N + row] += matA[col * N + scan] * matB[scan * N + row];
}
}
}
end = clock();


printf("A\n");

for (col = 0; col < N; col++){
for (row = 0; row < N; row++)
printf(" %3d",matA[col * N + row]);
printf("\n");
}
printf("\nB\n");

for (col = 0; col < N; col++){
for (row = 0; row < N; row++)
printf(" %3d",matB[col * N + row]);
printf("\n");
}
printf("\nC\n");

for (col = 0; col < N; col++){
for (row = 0; row < N; row++)
printf(" %6d",matC[col * N + row]);
printf("\n");
}
// 計算時間結果の表示
printf("It takes %d msec\n", end-start);

// 各行列のメモリ解放
free(matA);
free(matB);
free(matC);

return 0;
}

これをコマンドラインから cl c_matrix_3.c のあと実行。計算時間は0m秒。

matrix_3.cu:

#include <stdio.h>
#include <malloc.h>
#include <stdlib.h>
#include <time.h>

#define N 3 // 正方行列のサイズを指定(N×N)
#define BLOCK 1 // ブロックのサイズを指定

__global__ void
matrixMul(int* inMatA, int* inMatB, int* inMatC);

int main(int argc, char** ){

// 行列のサイズをバイト単位で算出
int matrixSize = sizeof(unsigned int) * N*N;

// ホスト側の行列変数設定
int* hMatA;
int* hMatB;
int* hMatC;
clock_t start, end;

// 行列変数のメモリ確保
hMatA = (int*)malloc(matrixSize);
hMatB = (int*)malloc(matrixSize);

// 初期値設定
int col, row;
for (col = 0; col < N; col++){
for (row = 0; row < N; row++){
hMatA[col * N + row] = rand() % 100;
hMatB[col * N + row] = rand() % 100;
}
}

// デバイス側の行列変数設定
int* dMatA;
int* dMatB;
int* dMatC;

// デバイスメモリ領域の確保
cudaMalloc((void**)&dMatA, matrixSize);
cudaMalloc((void**)&dMatB, matrixSize);
cudaMalloc((void**)&dMatC, matrixSize);

start = clock();
// ホストからデバイスへの変数の受け渡し
cudaMemcpy(dMatA, hMatA, matrixSize, cudaMemcpyHostToDevice);
cudaMemcpy(dMatB, hMatB, matrixSize, cudaMemcpyHostToDevice);

// ブロックサイズとグリッドサイズの設定
dim3 block(BLOCK, BLOCK);
dim3 grid( N / BLOCK, N / BLOCK);

// カーネルの起動
matrixMul<<<grid,block>>>(dMatA, dMatB, dMatC);
cudaThreadSynchronize();

// 結果の領域確保とデバイス側からのメモリ転送
hMatC = (int*)malloc(matrixSize);
cudaMemcpy(hMatC, dMatC, matrixSize, cudaMemcpyDeviceToHost);

end = clock();

printf("A\n");

for (col = 0; col < N; col++){
for (row = 0; row < N; row++)
printf(" %3d",hMatA[col * N + row]);
printf("\n");
}
printf("\nB\n");

for (col = 0; col < N; col++){
for (row = 0; row < N; row++)
printf(" %3d",hMatB[col * N + row]);
printf("\n");
}
printf("\nC\n");

for (col = 0; col < N; col++){
for (row = 0; row < N; row++)
printf(" %6d",hMatC[col * N + row]);
printf("\n");
}
// 計算時間結果の表示
printf("It takes %d msec\n", end-start);
// ホスト・デバイスメモリの解放
free(hMatA);
free(hMatB);
free(hMatC);
cudaFree(dMatA);
cudaFree(dMatB);
cudaFree(dMatC);

// 終了処理
cudaThreadExit();
}

__global__ void
matrixMul(int* inMatA, int* inMatB, int* inMatC){
int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
int scan;
int target = 0;

// 行列の演算を行う
for (scan = 0; scan < N; scan++) {
target += inMatA[col * N + scan] * inMatB[scan * N + row];
__syncthreads();
}

inMatC[col * N + row] = target;
}

これをコマンドラインから nvcc matrix_3.cu -o matrix_3 このあと実行。計算時間は0m秒。

よって、CPUの計算とGPUの並列計算はドロー。ちなみにCPU、オーバークロックをはたらかせ4.5GHzで計算させております。

次、1024行1024列の行列の掛け算。各々1行1列のみ表示させております。

cmatrix_1024.c:


#include <stdio.h>
#include <malloc.h>
#include <stdlib.h>
#include <time.h>

#define N 1024 // 正方行列のサイズを指定(N×N)

int main( int argc, char** argv)
{
int col, row, scan;

int* matA;
int* matB;
int* matC;
clock_t start, end;

// 行列A~Cのメモリ確保
matA = (int*)malloc(sizeof(int) * N*N);
matB = (int*)malloc(sizeof(int) * N*N);
matC = (int*)malloc(sizeof(int) * N*N);


// 行列A、Bの各要素に乱数を入れる
// 行列Cは初期値0
for (col = 0; col < N; col++){
for (row = 0; row < N; row++){
matA[col * N + row] = rand() % 100;
matB[col * N + row] = rand() % 100;
matC[col * N + row] = 0.0;
}
}

start = clock();

// 行列の計算C = A × B
for (col = 0; col < N; col++) {
for (row = 0; row < N; row++) {
for (scan = 0; scan < N; scan++) {
matC[col * N + row] += (float)(matA[col * N + scan] * matB[scan * N + row]);
}
}
}
end = clock();


printf("A\n");

for (col = 0; col < 1; col++){
for (row = 0; row < 1; row++)
printf(" %3d",matA[col * N + row]);
printf("\n");
}
printf("\nB\n");

for (col = 0; col < 1; col++){
for (row = 0; row < 1; row++)
printf(" %3d",matB[col * N + row]);
printf("\n");
}
printf("\nC\n");

for (col = 0; col < 1; col++){
for (row = 0; row < 1; row++)
printf(" %8d",matC[col * N + row]);
printf("\n");
}
// 計算時間結果の表示
printf("It takes %d msec\n", end-start);

// 各行列のメモリ解放
free(matA);
free(matB);
free(matC);

return 0;
}

コンパイルした後実行すると、5644m秒。

matrix_1024.cu:

#include <stdio.h>
#include <malloc.h>
#include <stdlib.h>
#include <time.h>

#define N 1024 // 正方行列のサイズを指定(N×N)
#define BLOCK 8 // ブロックのサイズを指定

__global__ void
matrixMul(int* inMatA, int* inMatB, int* inMatC);

int main(int argc, char** ){

// 行列のサイズをバイト単位で算出
int matrixSize = sizeof(unsigned int) * N*N;

// ホスト側の行列変数設定
int* hMatA;
int* hMatB;
int* hMatC;
clock_t start, end;

// 行列変数のメモリ確保
hMatA = (int*)malloc(matrixSize);
hMatB = (int*)malloc(matrixSize);

// 初期値設定
int col, row;
for (col = 0; col < N; col++){
for (row = 0; row < N; row++){
hMatA[col * N + row] = rand() % 100;
hMatB[col * N + row] = rand() % 100;
}
}

// デバイス側の行列変数設定
int* dMatA;
int* dMatB;
int* dMatC;

// デバイスメモリ領域の確保
cudaMalloc((void**)&dMatA, matrixSize);
cudaMalloc((void**)&dMatB, matrixSize);
cudaMalloc((void**)&dMatC, matrixSize);

start = clock();
// ホストからデバイスへの変数の受け渡し
cudaMemcpy(dMatA, hMatA, matrixSize, cudaMemcpyHostToDevice);
cudaMemcpy(dMatB, hMatB, matrixSize, cudaMemcpyHostToDevice);

// ブロックサイズとグリッドサイズの設定
dim3 block(BLOCK, BLOCK);
dim3 grid( N / BLOCK, N / BLOCK);

// カーネルの起動
matrixMul<<<grid,block>>>(dMatA, dMatB, dMatC);
cudaThreadSynchronize();

// 結果の領域確保とデバイス側からのメモリ転送
hMatC = (int*)malloc(matrixSize);
cudaMemcpy(hMatC, dMatC, matrixSize, cudaMemcpyDeviceToHost);

end = clock();

printf("A\n");

for (col = 0; col < 1; col++){
for (row = 0; row < 1; row++)
printf(" %3d",hMatA[col * N + row]);
printf("\n");
}
printf("\nB\n");

for (col = 0; col < 1; col++){
for (row = 0; row < 1; row++)
printf(" %3d",hMatB[col * N + row]);
printf("\n");
}
printf("\nC\n");

for (col = 0; col < 1; col++){
for (row = 0; row < 1; row++)
printf(" %8d",hMatC[col * N + row]);
printf("\n");
}
// 計算時間結果の表示
printf("It takes %d msec\n", end-start);
// ホスト・デバイスメモリの解放
free(hMatA);
free(hMatB);
free(hMatC);
cudaFree(dMatA);
cudaFree(dMatB);
cudaFree(dMatC);

// 終了処理
cudaThreadExit();
}

__global__ void
matrixMul(int* inMatA, int* inMatB, int* inMatC){
int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
int scan;
int target = 0;

// 行列の演算を行う
for (scan = 0; scan < N; scan++) {
target += inMatA[col * N + scan] * inMatB[scan * N + row];
__syncthreads();
}

inMatC[col * N + row] = target;
}

コンパイルして実行。計算時間はなんと脅威の93m秒。CPUの計算時間の1/56以下。CPU(Core i7 3770Kオーバークロック)より56倍速いです!!

ちなみにmatrix_1024.cuの変数BLOCKは16とか32とかにも設定してみてください。私のGTX670の場合、8が最速でした。

なんだか机上に乗るスパコンを持っている気分です。(^.^)




Palit GeForce GTX 670 JETSTREAM 2GB (NE5X670H1042-1042J V2)





Palit GeForce GTX 680 JETSTREAM 4GB (NE5X680010G2-1041J V2)


自作pc 俺ならばこれを買う。
スポンサーサイト

この記事へのコメント

コメントをお寄せ下さい

(コメント編集・削除に必要)

トラックバック

この記事のトラックバックURL
http://tj502.blog32.fc2.com/tb.php/1453-5a863dff

気になるコンピュータ情報

お勧め情報!

自作PCに関しては「自作コンピュータ 俺ならばこれを買う。」にまとめました。ご活用ください。



上記広告は1ヶ月以上更新のないブログに表示されています。新しい記事を書くことで広告を消せます。