■
なんかを書くのに調査。
nvccはアセンブリとして、謎のマシン語PTXを生成するのだが、なんかPTXのあとも色々最適化されてるように見えるな。
#include <stdio.h> __global__ void calc_clock(int *data, float *accum_store) { int tx = threadIdx.x; unsigned int b,e; float d = 0; int x=0, y=0, z=0, n=0, m=0, o=0; b = clock(); asm volatile ("" "add.s32 %1, %1, %1;\n\t" "add.s32 %2, %2, %2;\n\t" "add.s32 %3, %3, %3;\n\t" "add.s32 %4, %4, %4;\n\t" "add.s32 %5, %5, %5;\n\t" "add.s32 %6, %6, %6;\n\t" :"+f"(d), "+r"(x), "+r"(y), "+r"(z), "+r"(n), "+r"(m), "+r"(o) ); e = clock(); accum_store[0] = d + x + y + z + n + m + o; data[tx] = e-b; } int main() { int clock[512]; float accum[512]; int *d_clock; float *d_accum; cudaMalloc((void**)&d_clock, sizeof(int)*512); cudaMalloc((void**)&d_accum, sizeof(float)*512); calc_clock<<<1,1>>>(d_clock, d_accum); cudaMemcpy(clock, d_clock, sizeof(int)*512, cudaMemcpyDeviceToHost); cudaMemcpy(accum, d_accum, sizeof(float)*512, cudaMemcpyDeviceToHost); printf("%d %f\n", clock[0], accum[0]); }
例えば、こういうのは何の時間も計測できない。アセンブリは出力されているにも関わらず。
時間を計測するには、コンパイラには判定できない値をレジスタに入れておかないといけない
↓しかし、こうすると、今度は信じられないほどの時間がかかってしまう。
__global__ void calc_clock(int *data, float *accum_store) { int tx = threadIdx.x; unsigned int b,e; float d; int x=data[0], y, z, n, m, o; b = clock(); asm volatile ("add.s32 %1, %1, %1;\n\t" :"+f"(d), "+r"(x), "+r"(y), "+r"(z), "+r"(n), "+r"(m), "+r"(o) ); e = clock(); accum_store[0] = d + x + y + z + n + m + o; data[tx] = e-b; }
これはメモリのレイテンシはロードしたあと最初にレジスタに触ったときに発生するから。(これは普通のアーキテクチャと一緒)
これを解決するには、__syncthreads()を入れておくとよさげ。
__global__ void calc_clock(int *data, float *accum_store) { int tx = threadIdx.x; unsigned int b,e; float d; int x=data[0], y, z, n, m, o; __syncthreads(); // これね b = clock(); asm volatile ("add.s32 %1, %1, %1;\n\t" :"+f"(d), "+r"(x), "+r"(y), "+r"(z), "+r"(n), "+r"(m), "+r"(o) ); e = clock(); accum_store[0] = d + x + y + z + n + m + o; data[tx] = e-b; }
これで命令ごとのレイテンシを計測する準備ができた。やることは、
- レジスタにどこかのメモリから引っ張ってきた値を入れておく(値がいらないなら未初期値でも良いように見える)
- clock呼ぶ前に__syncthreads()
- 結果はどこかのメモリに書き込んでおく
だけどまだ問題があって、なんかちゃんとスケジューリングしてるぽくて。
asm volatile ("add.s32 %1, %1, %1;\n\t" "add.s32 %2, %2, %2;\n\t" "add.s32 %3, %3, %3;\n\t" "add.s32 %4, %4, %4;\n\t" "add.s32 %5, %5, %5;\n\t" "add.s32 %6, %6, %6;\n\t" "add.s32 %1, %1, %1;\n\t" :"+f"(d), "+r"(x), "+r"(y), "+r"(z), "+r"(n), "+r"(m), "+r"(o) );
asm volatile ("add.s32 %1, %1, %1;\n\t" "add.s32 %1, %1, %1;\n\t" "add.s32 %2, %2, %2;\n\t" "add.s32 %3, %3, %3;\n\t" "add.s32 %4, %4, %4;\n\t" "add.s32 %5, %5, %5;\n\t" "add.s32 %6, %6, %6;\n\t" :"+f"(d), "+r"(x), "+r"(y), "+r"(z), "+r"(n), "+r"(m), "+r"(o) );
このふたつの結果に違いが無い。
まあ、なんか、スケジューラの動作を予測しながらって感じで。
あとの問題として、インラインアセンブリは多分サポート外なのですぐにコンパイラが死ぬ。
上の点さえ押さえればメモリ操作以外はなんとか調べられるように見える。
asm volatile ("add.s32 %1, %1, %1;\n\t" // asm 外の命令のレイテンシから守ってくれるバリヤー "add.s32 %1, %1, %1;\n\t" "add.s32 %1, %1, %1;\n\t" "add.s32 %1, %1, %1;\n\t" "mov.s32 %7, %clock;\n\t" // 計測開始 "add.s32 %2, %2, %2;\n\t" // %1 は使わない "add.s32 %3, %3, %3;\n\t" "add.s32 %4, %4, %4;\n\t" "add.s32 %5, %5, %5;\n\t" "add.s32 %6, %6, %6;\n\t" "add.f32 %0, %0, %0;\n\t" "mov.s32 %8, %clock;\n\t" // 計測終了 "add.s32 %1, %1, %1;\n\t" "add.s32 %1, %1, %1;\n\t" "add.s32 %1, %1, %1;\n\t" "add.s32 %1, %1, %1;\n\t" "add.s32 %1, %1, %1;\n\t" "add.s32 %1, %1, %1;\n\t" // asm 外の命令のレイテンシから守ってくれるバリヤー
で、上で書いたようにレジスタに値突っこんで、最後に値使えばそれなりに正しい値が取れる気がする。
とりあえず、整数、浮動小数加算は18clockで、1warpが4cycleで実行されて、5warp(156thread)流しておけばレイテンシが隠蔽できるようなので、clockのクロックはシェーダのクロックと同じと見てよさげだな。
ただなんか、1warpだとスループット1/2になってるように見えるな。手動でレイテンシ隠蔽するのに慣れてる人でもなにをどうやっても2warp(64thread)流せということか。
GTX280とかは32SPあるので、つまり性能を出すには最低でも32x64 = 2048スレッドは作らないといけない。あとメモリがあるがさすがに時間やばいので明日。
ららびんは16core x 16SIMD = 256スレッドか。あとGTXは次世代で倍ぐらいになるらしいので4000スレッドか。まあ、2000スレッド分も並列性があるようなプログラムは多分4000ぐらい並列性はあるだろうからもう一緒だろという気はするな。
並列性のスケール感は、
- 1 → 依存が多いアプリ
- 2 → タスクが複数あるアプリ
- 10 → なんかプログラマが工夫してタスクをうまく分割したアプリ
- 1000000 → なんか
のどれかという感じ。
