初めてのCUDA

ここまでのあらすじ。9800GTを積んだマシンでKNOPPIX for CUDAを起動した。KNOPPIX for CUDAが"Can't find KNOPPIX filesystem"で起動できない問題(解決)

KNOPPIXあまり詳しくないので適当にいじってみる。CDからブートしているからもしかしてファイルに書き込めないかなと思ってホームにファイルを作ったらちゃんと書ける。RAMディスクなのかな、と思ったらホームは/ramdisk/home/knoppix/だった。つまり電源を切れば消えるんだな。USBメモリを指してみたらちゃんと認識したので消えて困るものはここに入れればいいな。

/usr/local/cudaに色々入っている。とりあえず/usr/local/cuda/NVIDIA_CUDA_SDK/projectsに入っているたくさんのサンプルの中から適当に初歩そうなものをホームにコピーしてみよう。「clock」を選んだ。cuda/bin/nvccでclock.cuをコンパイルしようとしたらcutil.hがないと怒られる。っていうかMakefileあるじゃん、これを読もう。

Makefileの中では「include ../../common/common.mk」とかやっているからとりあえず絶対パスに書き換える。(注:最終的にこれは必要なかった)
emacsは入っていない。なんかkwriteってのが入っているからそれをとりあえず使ってみた。あ、vimの7.0が入っている。これを使えばいいな。

makeしてみた。「../../libを作成できません」ふむ、そうか。CUDA_SDKディレクトリ全体を書き込み可能なところにコピーするのが手っ取り早かったな。そうしよう。nvccにパスも通しておく。

makeしてもnvcc -I ../../common/inc clock.cuしても下のようなエラーが出る。

"/usr/include/c++/4.2.1/i586-suse-linux/bits/c++config.h", line 149: error:
         expected a "{"
 namespace std __attribute__ ((__visibility__ ("default"))) {

NVIDIA Forums > nvcc problem with __attribute__によればこれはgcc4.2での変更が原因で、gcc4.1を使えばいいとのこと。幸いKNOPPIX for CUDAにはgcc4.1も入っていた。フォーラムに書いてあるように適当な位置にシンボリックリンクを作って、common/common.mkの「# Compilers」の所のnvccに--compilar-bindirを付けた。あとついでにgccやg++もgcc-4.1やg++-4.1に変更した。たぶん必要だろうと思ったので。

makeが成功するようになった。CUDA_SDK/bin/linux/releaseに実行可能なバイナリが出来ている。

• • • • -

とりあえず単純に9 * 9の掛け算の九九の表を並列に実行してみたつもり。IS2Dの方が速かった。あたりまえだけど。カーネルの名前がtimedReductionなのはもとにしたサンプルのままだなー。ところで本当は-DIS2Dとかで挙動を切り替えて試したかったのだけどmake経由でコンパイルしてたからダメだった。他の人はどうしてるのかな。

// calc multiplication table
// derived from CUDA_SDK/project/clock

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

#include <cutil.h>
#define N 9
#define NUM_BLOCKS 1
#define IS2D

__global__ static void timedReduction(const int* input, int* output, clock_t * timer) {
#ifdef IS2D
    const int x = threadIdx.x, y = threadIdx.y, tid = x * y;
#else
    const int tid = threadIdx.x, x = tid % N, y = tid / N;
#endif

    const int bid = blockIdx.x;
    if (tid == 0) timer[bid] = clock();
    output[tid] = input[x] * input[y];
    __syncthreads();
    if (tid == 0) timer[bid+gridDim.x] = clock();
}

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

#ifdef IS2D
    dim3 threads(N, N);
#else
    dim3 threads(N * N);
#endif

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

    clock_t timer[NUM_BLOCKS * 2];
    int input[N];
    int output[N * N];

    for (int i = 0; i < N; i++)
    {
        input[i] = i + 1;
    }

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

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

    timedReduction<<<NUM_BLOCKS, threads>>>(dinput, doutput, dtimer);

    CUDA_SAFE_CALL(cudaMemcpy(output, doutput, sizeof(int) * N * N, 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));

    // Compute the difference between the last block end and the first block start.
    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;
    }

    // print result
    for(int i=0; i < N; i++){
    	for(int j=0; j < N; j++){
	    printf("%3d, ", output[i + j * N]);
	}
	printf("\n");
    }
    printf("time = %d\n", maxEnd - minStart);

    CUT_EXIT(argc, argv);
}