Powered by Twitter Tools.
Archive for 7月 20th, 2007
実はAdvanced/W-Zero3[es]が欲しいかもしれない
とりあえずスマートフォンが欲しいのですよ。
で、一応Softbankからも出てるし、来月にはX01Tが出るわけなんだが。
どうもサイトを見ると料金プランがわからん。パケットし放題に「月額1,029円～4,410円で、国内はメールもウェブもし放題！（ＰＣサイトブラウザ除く）」とか書いてあるし。
っつーわけで、圏外率最強クラスの3Gを980円プランに変更して（変更しても家族間通話とかは安いはずだし）、アドエス買っちゃうのが一番幸せになれるかなあとか考えているわけだ。
どうせプリンタとかレコーダとかいじるために週末に秋葉ヨドバシあたりへ行ってこようと思っているので、ついでに料金のこととかも聞いてくるか。
ほとんど通話をしないから料金なんて対してかからないぜふふーんとか思っていると、たまにサポセンだのなんだのにかけた後に料金表示を見て３桁の数字が出てくるのが納得いかないんだよなあ。
CUDAでN-QUEENやってみた
主に、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);
- }