CUDA日記1 - 西尾泰和のはてなダイアリー

スレッドのsyncはブロック単位。ブロックごとにあるshared memoryはレジスタ並みに高速。協調する必要があるスレッドは同じブロックにいるべき。

GPU Gemsを見てGPUPUではscatter(ランダムアクセス書き込み)が困難だからアルゴリズムを工夫する必要がある、と認識していたがCUDAの長所はそれを可能にしたところらしい。

カーネルの実行は非同期だが、結果を取り出すMemcpyがカーネルの実行を待つので同期処理しているように見えていたみたいだ。

ホストからいじれるのは__device__のメモリだけ。__device__はレイテンシが大きい。__shared__の数百倍。テーブル引きをするような場合、__device__に入れっぱなし二するのは愚の骨頂なので__shared__にコピーすべきだろうか。それともカーネルの中に埋め込んだ方がいいのかな？もちろんカーネルのサイズに限界はあるのだろうけど。テーブルのサイズとアクセスの量に依存するかな。

カーネルの中で何も修飾子を書かずに宣言している変数は__shared__になる。ただしサイズが5以上の配列は__device__になる。

int4とかの「ベクトル型」がある。これって、例えば足し算したら並列に足されるのかな？後で試そう。

cudaMemcpyAsyncなんてのがあるな。

- - - -

そんなこんなで最中限の一番最後のターンで最善手が何かを全探索するコードを書いてみた。とりあえずエラーなく動くところまでやってみただけなので結果が正しいかどうかはまだ検証していないけど、実行するとtime = 10232で-1, 23, 3という結果が帰ってくる。僕がプレイするときでもこの曲面では真ん中の5を出すので、あってるのかな〜。速度はブロックも1個しか使っていないし問題が小さすぎて面白みがないから最終ラウンド全体の読み切りまでやるようにしてからC++でのコードと比較しよう。

あ、とりあえずx == yのケースを省くの忘れてるな。

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

#include <cutil.h>
#define NUM_BLOCKS    1

__device__ static int calc_median(int v0, int v1, int v2){
    if(v1 < v2){
    	if(v2 < v0){
	    return v2;
	}else if(v0 < v1){
	    return v1;
	}else{
	    return v0;
	}
    }else{
    	if(v1 < v0){
	    return v1;
	}else if(v0 < v2){
	    return v2;
	}else{
	    return v0;
	}
    }
}

__global__ static void kernel(const int * input, int * output, clock_t * timer)
{
    __shared__ int shared[3 * 7 * 7];

    const int x = threadIdx.x, y = threadIdx.y, z = threadIdx.z;
    const int bid = blockIdx.x;
    const int tid = x + y * 7 + z * 7 * 7;
    if (tid == 0) timer[bid] = clock();
    
    // update round score
    int buf[] = {input[10], input[11], input[12]};
    int my=input[z], o1=input[3 + x], o2=input[3 + y];
    int median = calc_median(my, o1, o2);
    if(my == median) buf[0] += median / 4 + 1;
    if(o1 == median) buf[1] += median / 4 + 1;
    if(o2 == median) buf[2] += median / 4 + 1;
    
    // update game score
    median = calc_median(buf[0], buf[1], buf[2]);
    o1 = buf[1]; o2 = buf[2]; my = buf[0];
    buf[0] = input[13]; // game score
    buf[1] = input[14];
    buf[2] = input[15];
    if(my == median) buf[0] += median;
    if(o1 == median) buf[1] += median;
    if(o2 == median) buf[2] += median;

    // find winner
    o1 = buf[1]; o2 = buf[2]; my = buf[0];
    median = calc_median(buf[0], buf[1], buf[2]);
    int win = 0;
    if(my == median) win += 2;
    if(o1 == median) win--;
    if(o2 == median) win--;
    
    shared[tid] = win;

    __syncthreads();
    if(x == 0){
        int sum = 0;
    	for(int i=0; i < 7; i++){
	    sum += shared[tid + 7 * i];
	}
	shared[tid] = sum;
    }
    __syncthreads();
    if(x == 0 && y == 0){
        int sum = 0;
    	for(int i=0; i < 7; i++){
	    sum += shared[tid + i];
	}
	output[z] = sum;
    }
    __syncthreads();
    if (tid == 0) timer[bid+gridDim.x] = clock();
}

int main(int argc, char** argv)
{
    CUT_DEVICE_INIT();

    int * dinput = NULL;
    int * doutput = NULL;
    clock_t * dtimer = NULL;

    clock_t timer[NUM_BLOCKS * 2];
    int input[] = {
    	1, 5, 13, // own cards
    	4, 15, 6, 7, 8, 9, 10, // others cards
	0, 5, 0, // round score
	0, 5, 0}; // game score

    CUDA_SAFE_CALL(cudaMalloc((void**)&dinput, sizeof(int) * 16));
    CUDA_SAFE_CALL(cudaMalloc((void**)&doutput, sizeof(int) * 3));
    CUDA_SAFE_CALL(cudaMalloc((void**)&dtimer, sizeof(clock_t) * NUM_BLOCKS * 2));

    CUDA_SAFE_CALL(cudaMemcpy(dinput, input, sizeof(int) * 16, cudaMemcpyHostToDevice));

    dim3 THREADS(7, 7, 3);
    kernel<<<NUM_BLOCKS, THREADS>>>(dinput, doutput, dtimer);

    int output[3];
    CUDA_SAFE_CALL(cudaMemcpy(output, doutput, sizeof(int) * 3, cudaMemcpyDeviceToHost));
    CUDA_SAFE_CALL(cudaMemcpy(timer, dtimer, sizeof(clock_t) * NUM_BLOCKS * 2, cudaMemcpyDeviceToHost));

    CUDA_SAFE_CALL(cudaFree(dinput));
    CUDA_SAFE_CALL(cudaFree(doutput));
    CUDA_SAFE_CALL(cudaFree(dtimer));

    clock_t minStart = timer[0];
    clock_t maxEnd = timer[NUM_BLOCKS];

    for (int i = 1; i < NUM_BLOCKS; i++)
    {
        minStart = timer[i] < minStart ? timer[i] : minStart;
        maxEnd = timer[NUM_BLOCKS+i] > maxEnd ? timer[NUM_BLOCKS+i] : maxEnd;
    }

    printf("time = %d\n", maxEnd - minStart);
    printf("%d, %d, %d\n", output[0], output[1], output[2]);
    CUT_EXIT(argc, argv);
}