GPGPU#2

常駐というより、巡回しているだけだけどね。
折角だから、nvccの出力(ptx)でも張ってみようか。先ずはソース。
--
#include <cstdio> // WIN版だとリンクにvcを使う所為かiostreamを巧く使えない
const int NofThreads = 256; // 32以下なら1warp内で完結、最大512
const int NofBlocks = 256; // 物理ブロック数をオーバーしたら時分割処理される
// device側コード
#include <cuda_runtime.h>
__device__ float dVals[NofBlocks * NofThreads]; // >45も書いている、device側のメモリ
__global__ void foo(float ratio)
{
float vs, vc;
unsigned offset = blockIdx.x * blockDim.x + threadIdx.x; // この右辺の変数で何番目のデータスレッドか判る
__sincosf(offset * ratio, & vs, & vc); // sincosf()の組み込み命令版
dVals[offset] = vs * vs + vc * vc;
}
// host(CPU)側コード
const float M_PI = 3.14159265358979323846f;
int main()
{
foo<<<NofBlocks, NofThreads>>>(M_PI / (NofBlocks * NofThreads)); // ここで並列起動
float hVals[NofBlocks * NofThreads]; // host(CPU)側メモリ
// この場合はcudaMemcpyFromSymbol()を使用(cudaMalloc()した場合はcudaMemcpy())
cudaMemcpyFromSymbol(hVals, dVals, sizeof(* hVals) * NofBlocks * NofThreads);
for (unsigned ic = 0; ic < NofBlocks * NofThreads; ++ic) {
printf("%g ", hVals[ic]);
}
fputchar('\n');
return 0;
}
// 所要時間はCore2Duo 6320 @ 1.86GHzで1秒掛からない程度
// シンプルすぎて時間測定には不向きなので要注意

んで、こっちがnvcc -ptx -O3した出力の抜粋
--
.global .align 4 .b8 dVals[262144];

.entry foo
{
.reg .u32 $r1,$r2,$r3,$r4,$r5,$r6,$r7,$r8;
.reg .f32 $f1,$f2,$f3,$f4,$f5,$f6,$f7;
.param .f32 __cudaparm_ratio;
.loc 12 10 0
$LBB1_foo:
.loc 12 15 0
cvt.u32.u16 $r1, %ctaid.x; //
cvt.u32.u16 $r2, %ntid.x; //
mul.lo.u32 $r3, $r1, $r2; //
cvt.u32.u16 $r4, %tid.x; //
add.u32 $r5, $r4, $r3; //
cvt.rn.f32.u32 $f1, $r5; //
ld.param.f32 $f2, [__cudaparm_ratio]; // id:31 __cudaparm_ratio+0x0
mul.f32 $f3, $f2, $f1; //
cos.f32 $f4, $f3; //
sin.f32 $f5, $f3; //
mul.f32 $f6, $f5, $f5; //
mad.f32 $f7, $f4, $f4, $f6; //
mov.u32 $r6, dVals; //
mul.lo.u32 $r7, $r5, 4; //
add.u32 $r8, $r6, $r7; //
st.global.f32 [$r8+0], $f7; // id:32 dVals+0x0
exit; //
} // foo
--
# 解説要るなら言っとくれ