主に、CUDAでプログラムを組むとこんな感じだよーってのをチェックするため。こんなもの高速化しても論文になりませ……ならないと思います。多分。
とりあえず問題サイズと同じだけの並列化*もやってちょうど200行くらい。少なくともシェーダを叩くよりはやりやすいと思う。メモリの扱いだけ気をつければ、単純なCで書ける。まあそのメモリの扱いがめんどいんだけどさ。
（* : N=8の時は8並列。要するに、IDがｎの人は１列目のクイーンをｎ行目に置くと仮定してネ、っていう並列化。）
問題サイズを12より大きくするとこけるのはなんだろう。とりあえずでっちあげただけなのでちゃんと調査してません。もしかしたらスレッド（CUDAの実行モデルで言うところの一番小さい単位。スレッドを束にしたブロックとかいう単位もある。まぁいわゆるスレッド並列化のスレッドと考えていいよ。）の何らかの制限を突破してるかもね。メモリ量とか時間とか。シェーダみたいにループの回数制限とかあったっけかなぁ？
ちなみに実行時間を測定したところ、CPUの非並列非再帰版より明らかに遅かった。更に細かく並列化して高速化する余地というかアイディアはありまくりなので、この後もうちょっと頑張って高速化するかもね。しないかもね。
一応フルソース載せておきます。参考になるかは知らない。とっとと帰ってゼミ資料でも捏造して寝ます。CUDAたんと戯れていたら東急ストアの特売日逃したよチクショウ！
CUDA:
- #include <stdlib.h>
- #include <stdio.h>
- #include <string.h>
- #include <math.h>
- #include <unistd.h>
- #include <cutil.h>
- #define N 12
- __global__ void test(
- int size,
- int* numAnswers,
- int* E,
- int* SE,
- int* NW,
- int* answer,
- int* etc
- )
- {
- int id = threadIdx.x;
- if(id>size){
- if(id<size*2){
- etc[id] = -1;
- }
- return;
- }
- etc[id] = 1;
- if(threadIdx.y>0){
- return;
- }
- int offset1 = id*size;
- int offset2 = id*(size*2-1);
- int x, y;
- x=0;
- y=id;
- answer[offset1+x] = y;
- E[offset1+y] = 1;
- SE[offset2+(size-1)-x+y] = 1;
- NW[offset2+x+y] = 1;
- x++;
- y=0;
- while(1){
- if(x==size){
- // all OK
- //printAnswer(d);
- numAnswers[id]++;
- x--;
- y = answer[offset1+x];
- //removeQueen(d,x,y);
- E[offset1+y] = 0;
- SE[offset2+(size-1)-x+y] = 0;
- NW[offset2+x+y] = 0;
- y++;
- }else if(y==size){
- x--;
- if(x<1){
- break;
- }
- y = answer[offset1+x];
- //removeQueen(d,x,y);
- E[offset1+y] = 0;
- SE[offset2+(size-1)-x+y] = 0;
- NW[offset2+x+y] = 0;
- y++;
- }else{
- int test=0;
- if(E[offset1+y] || SE[offset2+(size-1)-x+y] || NW[offset2+x+y]){
- test=0;
- }else{
- answer[offset1+x] = y;
- E[offset1+y] = 1;
- SE[offset2+(size-1)-x+y] = 1;
- NW[offset2+x+y] = 1;
- test=1;
- }
- if(test==1){
- x++;
- y=0;
- }else{
- y++;
- }
- }
- }
- //etc[id] = offset;
- }
- void
- runTest(int argc, char** argv)
- {
- CUT_DEVICE_INIT();
- printf("%d Byte\n", (1+N+N*N+(N*(N*2-1))+(N*(N*2-1))+N*N)*4);
- // CPU側のデータを準備
- int size;
- int numAnswers[N];
- int E[N*N];
- int SE[N*(N*2-1)];
- int NW[N*(N*2-1)];
- int answer[N*N];
- int etc[N*2];
- int i;
- size = N;
- for(i=0; i<N; i++){
- numAnswers[i] = 0;
- }
- for(i=0; i<N*2; i++){
- etc[i] = 0;
- }
- for(i=0; i<N*N; i++){
- E[i] = 0;
- answer[i] = 0;
- }
- for(i=0; i<N*(N*2-1); i++){
- SE[i] = 0;
- NW[i] = 0;
- }
- // GPU側のメモリを準備
- int* d_numAnswers;
- int* d_E;
- int* d_SE;
- int* d_NW;
- int* d_answer;
- int* d_etc;
- CUDA_SAFE_CALL(cudaMalloc((void**)&d_numAnswers, sizeof(int)*N));
- CUDA_SAFE_CALL(cudaMalloc((void**)&d_E, sizeof(int)*N*N));
- CUDA_SAFE_CALL(cudaMalloc((void**)&d_SE, sizeof(int)*(N*2-1)*N));
- CUDA_SAFE_CALL(cudaMalloc((void**)&d_NW, sizeof(int)*(N*2-1)*N));
- CUDA_SAFE_CALL(cudaMalloc((void**)&d_answer, sizeof(int)*N*N));
- CUDA_SAFE_CALL(cudaMalloc((void**)&d_etc, sizeof(int)*N*2));
- CUDA_SAFE_CALL(cudaMemcpy(d_numAnswers, &numAnswers, sizeof(int)*N, cudaMemcpyHostToDevice) );
- CUDA_SAFE_CALL(cudaMemcpy(d_E, E, sizeof(int)*N*N, cudaMemcpyHostToDevice) );
- CUDA_SAFE_CALL(cudaMemcpy(d_SE, SE, sizeof(int)*(N*2-1)*N, cudaMemcpyHostToDevice) );
- CUDA_SAFE_CALL(cudaMemcpy(d_NW, NW, sizeof(int)*(N*2-1)*N, cudaMemcpyHostToDevice) );
- CUDA_SAFE_CALL(cudaMemcpy(d_answer, &answer, sizeof(int)*N*N, cudaMemcpyHostToDevice) );
- CUDA_SAFE_CALL(cudaMemcpy(d_etc, &etc, sizeof(int)*N*2, cudaMemcpyHostToDevice) );
- // 実行のためのパラメタを準備
- dim3 threads(N, 1, 1);
- dim3 grid(1,1,1);
- // 実行
- unsigned int hTimer;
- CUT_SAFE_CALL( cutCreateTimer(&hTimer) );
- CUT_SAFE_CALL( cutResetTimer(hTimer) );
- CUT_SAFE_CALL( cutStartTimer(hTimer) );
- test<<<grid, threads>>>(size, d_numAnswers, d_E, d_SE, d_NW, d_answer, d_etc);
- CUDA_SAFE_CALL( cudaThreadSynchronize() );
- CUT_SAFE_CALL( cutStopTimer(hTimer) );
- double gpuTime = cutGetTimerValue(hTimer);
- printf("Time: %f ms\n", gpuTime);
- // GPUの処理に問題が起きていないかの確認
- CUT_CHECK_ERROR("Kernel execution failed");
- // 演算結果の取得
- CUDA_SAFE_CALL(cudaMemcpy(&numAnswers, d_numAnswers, sizeof(int)*N, cudaMemcpyDeviceToHost) );
- CUDA_SAFE_CALL(cudaMemcpy(&etc, d_etc, sizeof(int)*N, cudaMemcpyDeviceToHost) );
- // 演算結果の確認
- int nAll=0;
- for(i=0; i<N*2; i++){
- printf(" %d", etc[i]);
- }
- printf("\n");
- for(i=0; i<N; i++){
- nAll += numAnswers[i];
- printf(" %d\n", numAnswers[i]);
- }
- printf("sum %d\n", nAll);
- // クリーンアップ
- //free();
- CUDA_SAFE_CALL(cudaFree(d_numAnswers));
- CUDA_SAFE_CALL(cudaFree(d_E));
- CUDA_SAFE_CALL(cudaFree(d_SE));
- CUDA_SAFE_CALL(cudaFree(d_NW));
- CUDA_SAFE_CALL(cudaFree(d_answer));
- CUDA_SAFE_CALL(cudaFree(d_etc));
- }
- int
- main(int argc, char** argv)
- {
- runTest(argc, argv);
- CUT_EXIT(argc, argv);
- }