CUDA 5の新機能(5): Texture Cacheを自動的に使うDevice Memory読み込み

出典: トータル・ディスクロージャ・サイト(事実をありのままに)

「CUDA 5の新機能」シリーズの筆を休めて先日記述した、Fortranのサブルーチン引数のAliasingの記事。実はこの記事の内容は、CUDA 5で対応したTesla K20向けのCompute Capablity 3.5で、Device Memory読み込みでTexture Cacheを自動的に使うコードを出力するための条件に密接に関わっている。

__restrict__修飾子

C99のrestrict修飾子と同等の機能を持つGCC拡張に、__restrict__という修飾子が在る。restrictと使い方は同じであるが、GCC拡張を使わないコンパイルオプションを指定しない限り、C言語のどの規格を指定しても有効であるうえ、C++でも同等の機能を使うことができる(C++がC言語から分岐した後に制定されたC99はC++の規格とは独立していて、C99の機能の中には最新のC++の規格でも追随していないものが有り、restrictもそういう機能の一つ)。

このGCC拡張はCUDAも取り入れていて、たとえば次のようなCUDAプログラムが在る。

__global__ void
f0(float *a, float *b)
{
  a[threadIdx.x] = 1.0f;
  a[threadIdx.x] = b[threadIdx.x];
}

__global__ void
f1(float *__restrict__ a, float *__restrict__ b)
{
  a[threadIdx.x] = 1.0f;
  a[threadIdx.x] = b[threadIdx.x];
}

これをコンパイルすると、次のようなPTXが出力される。

	.entry _Z2f0PfS_ (
		.param .u64 __cudaparm__Z2f0PfS__a,
		.param .u64 __cudaparm__Z2f0PfS__b)
	{
	.reg .u64 %rd<8>;
	.reg .f32 %f<4>;
	.loc	15	2	0
$LDWbegin__Z2f0PfS_:
	.loc	15	4	0
	cvt.u64.u16 	%rd1, %tid.x;
	mul.lo.u64 	%rd2, %rd1, 4;
	ld.param.u64 	%rd3, [__cudaparm__Z2f0PfS__a];
	add.u64 	%rd4, %rd3, %rd2;
	mov.f32 	%f1, 0f3f800000;     	// 1
	st.global.f32 	[%rd4+0], %f1;
	.loc	15	5	0
	ld.param.u64 	%rd5, [__cudaparm__Z2f0PfS__b];
	add.u64 	%rd6, %rd5, %rd2;
	ld.global.f32 	%f2, [%rd6+0];
	st.global.f32 	[%rd4+0], %f2;
	.loc	15	6	0
	exit;
$LDWend__Z2f0PfS_:
	} // _Z2f0PfS_

	.entry _Z2f1PfS_ (
		.param .u64 __cudaparm__Z2f1PfS__a,
		.param .u64 __cudaparm__Z2f1PfS__b)
	{
	.reg .u64 %rd<8>;
	.reg .f32 %f<3>;
	.loc	15	9	0
$LDWbegin__Z2f1PfS_:
	.loc	15	12	0
	cvt.u64.u16 	%rd1, %tid.x;
	mul.lo.u64 	%rd2, %rd1, 4;
	ld.param.u64 	%rd3, [__cudaparm__Z2f1PfS__b];
	add.u64 	%rd4, %rd3, %rd2;
	ld.global.f32 	%f1, [%rd4+0];
	ld.param.u64 	%rd5, [__cudaparm__Z2f1PfS__a];
	add.u64 	%rd6, %rd5, %rd2;
	st.global.f32 	[%rd6+0], %f1;
	.loc	15	13	0
	exit;
$LDWend__Z2f1PfS_:
	} // _Z2f1PfS_

f0のプログラムは

  1. aに1.0fを書き込む(st.global.f32)
  2. bの値をレジスタに読み込む(ld.global.f32)
  3. aにレジスタの値を書き込む(st.global.f32)

の3つの部分から成っているのに対して、f1のプログラムは

  1. bの値をレジスタに読み込む(ld.global.f32)
  2. aにレジスタの値を書き込む(st.global.f32)

の2つの部分だけと成っている。最初のaに1.0fを書き込む処理は、変数b__restrict__修飾子が付いているf1では、bの値に影響を及ぼさないとみなして無駄な処理だとして最適化で省略されるのである。

Compute Capability 3.5での__restrict__修飾子の扱い

先ほどのプログラムを、CUDA 5のNVCCで-arch=sm_35のコンパイルオプションを付けてコンパイルする。

.visible .entry _Z2f0PfS_(
	.param .u64 _Z2f0PfS__param_0,
	.param .u64 _Z2f0PfS__param_1
)
{
	.reg .s32 	%r<6>;
	.reg .f32 	%f<2>;
	.reg .s64 	%rd<8>;


	ld.param.u64 	%rd1, [_Z2f0PfS__param_0];
	ld.param.u64 	%rd2, [_Z2f0PfS__param_1];
	cvta.to.global.u64 	%rd3, %rd2;
	.loc 3 4 1
	mov.u32 	%r1, %tid.x;
	cvta.to.global.u64 	%rd4, %rd1;
	.loc 3 4 1
	mul.wide.u32 	%rd5, %r1, 4;
	add.s64 	%rd6, %rd4, %rd5;
	mov.u32 	%r2, 1065353216;
	.loc 3 4 1
	st.global.u32 	[%rd6], %r2;
	.loc 3 5 1
	add.s64 	%rd7, %rd3, %rd5;
	ld.global.f32 	%f1, [%rd7];
	st.global.f32 	[%rd6], %f1;
	.loc 3 6 2
	ret;
}

.visible .entry _Z2f1PfS_(
	.param .u64 _Z2f1PfS__param_0,
	.param .u64 _Z2f1PfS__param_1
)
{
	.reg .s32 	%r<3>;
	.reg .f32 	%f<2>;
	.reg .s64 	%rd<8>;


	ld.param.u64 	%rd1, [_Z2f1PfS__param_0];
	ld.param.u64 	%rd2, [_Z2f1PfS__param_1];
	cvta.to.global.u64 	%rd3, %rd2;
	.loc 3 11 1
	mov.u32 	%r1, %tid.x;
	cvta.to.global.u64 	%rd4, %rd1;
	.loc 3 11 1
	mul.wide.u32 	%rd5, %r1, 4;
	add.s64 	%rd6, %rd4, %rd5;
	.loc 3 12 1
	add.s64 	%rd7, %rd3, %rd5;
	ld.global.nc.f32 	%f1, [%rd7];
	st.global.f32 	[%rd6], %f1;
	.loc 3 13 2
	ret;
}

f1に在ったld.global.f32命令が、ld.global.nc.f32命令に置き換わっている。ld.global命令のncという接尾辞は「non-coherent cache」の略であり、データ書き込みに対するデータの一貫性を持たないキャッシュ、つまりTexture Cacheを自動的に使うことを示している。__restrict__修飾子により、bの参照先のデータの値は他のポインタの参照先への書き込みで変化することは無いというヒントが与えられているため、このような最適化が可能となっている。

Tesla K20の現物は当然ながら手元に無いため、残念ながら速度比較はできないが、わざわざこういう新命令を作ってくる以上性能向上は有るとみて良いだろう。どのくらいの性能向上となるか、現物を手に入れ次第測定してみたいところである。


この記事へのコメントをお寄せください

  • サイトへの書き込みに差し支えございましたら トータルディスクロージャーサイトサポート係へメールをお送りください
  • トータル・ディスクロージャ・サイトに投稿された文章と画像は、すべてその著作権がHPCシステムズ株式会社に帰属し、HPCシステムズ株式会社が著作権を所有することに同意してください。
  • あなたの文章が他人によって自由に編集、配布されることを望まない場合は、投稿を控えてください。
  • コメントを書き込む場合は名前にひらがなを織り交ぜてください。
  • あなたの投稿する文章と画像はあなた自身によって書かれたものであるか、パブリック・ドメインかそれに類する自由なリソースからの複製であることを約束してください。あなたが著作権を保持していない作品を許諾なしに投稿してはいけません!

<comments hideform="false" />


Comments

ノート:CUDA 5の新機能(5): Texture Cacheを自動的に使うDevice Memory読み込み

個人用ツール