概要

SASS とは何か？SASSについて知って何の役に立つのかについて。

SASS とは何か

PTXのさらに後ろにあるほぼネイティブアセンブリ。

PTX を ptxas に入れると、NVIDIA GPU用の機械語を含むcubinが出るが、 この cubin をcuobjdumpを使って逆アセンブルした結果として確認できる。

asfermi という非公式のツールを使えば SASS->cubin へのアセンブルもできる

SASS が何の略かは公式にはわからないが、Shader ASSembly ではないかとの説が。

SASS の歴史(というほど大したものではないが)

cuobjdump 以前

人が頑張って機械語を解析して、独自に作った decuda というツールがあった。

このdecudaを活用して書かれた "Micro-benchmarking the GT200" は全CUDAプログラマが読むべきものであると思う。

cuobjdump 以降

(多分)CUDA3.0になって、cuobjdumpが公式に入った。これによってdecudaは必要無くなった。

SASS の読みかた

1. nvcc を使って cubinを出す

 $ cat a.cu
__global__ void func(int *p) {
    p[0]++;
}
 $ nvcc -cubin a.cu
 

2. cuobjdump を使って SASS を見てニヤニヤする

 $ cuobjdump -sass a.cubin

	code for sm_10
		Function : _Z4funcPi
	/*0000*/     /*0x1000c8010423c780*/ 	MOV R0, g [0x4];
	/*0008*/     /*0xd00e000580c00780*/ 	GLD.U32 R1, global14 [R0];
	/*0010*/     /*0x2001820500000003*/ 	IADD32I R1, R1, 0x1;
	/*0018*/     /*0xd00e0005a0c00781*/ 	GST.U32 global14 [R0], R1;
		..........................
 

cubin はただのELFなので、readelfでもシンボルの確認等ができる

 $ nm a.cubin
00000000 T _Z4funcPi
 $ readelf -h a.cubin
ELF Header:
  Magic:   7f 45 4c 46 01 01 01 33 02 00 00 00 00 00 00 00 
  Class:                             ELF32
  Data:                              2's complement, little endian
  Version:                           1 (current)
  OS/ABI:                            <unknown: 33>
  ABI Version:                       2
  Type:                              EXEC (Executable file)
  Machine:                           NVIDIA CUDA architecture
  Version:                           0x1
  Entry point address:               0x0
  Start of program headers:          668 (bytes into file)
  Start of section headers:          52 (bytes into file)
  Flags:                             0xa010a
  Size of this header:               52 (bytes)
  Size of program headers:           32 (bytes)
  Number of program headers:         3
  Size of section headers:           40 (bytes)
  Number of section headers:         7
  Section header string table index: 1

arch

-arch を変えると結構内容が変わる

 $ nvcc -cubin -arch sm_20 a.cu
 $ cubojdump -sass a.cubin
	code for sm_20
		Function : _Z4funcPi
	/*0000*/     /*0x00005de428004404*/ 	MOV R1, c [0x1] [0x100];
	/*0008*/     /*0x80009de428004000*/ 	MOV R2, c [0x0] [0x20];
	/*0010*/     /*0x00201c8588000000*/ 	LDU R0, [R2];
	/*0018*/     /*0x04001c034800c000*/ 	IADD R0, R0, 0x1;
	/*0020*/     /*0x00201c8590000000*/ 	ST [R2], R0;
	/*0028*/     /*0x00001de780000000*/ 	EXIT;
		..........................
 $ nvcc -cubin -arch sm_30 a.cu
 $ cuobjdump -sass a.cubin 

	code for sm_30
		Function : _Z4funcPi
	/*0008*/     /*0x10005de428004001*/ 	MOV R1, c [0x0] [0x44];
	/*0010*/     /*0x00009de428004005*/ 	MOV R2, c [0x0] [0x140];
	/*0018*/     /*0x00201c8580000000*/ 	LD R0, [R2];
	/*0020*/     /*0x04001c034800c000*/ 	IADD R0, R0, 0x1;
	/*0028*/     /*0x00201c8590000000*/ 	ST [R2], R0;
	/*0030*/     /*0x00001de780000000*/ 	EXIT;
	/*0038*/     /*0xe0001de74003ffff*/ 	BRA 0x38;
		..........................
 

arch は、nvcc.pdf に詳しく書いてあるが、何かの役に立つと思うのでいくらか解説しておく。

arch は、使える機能を決める。code は、cubinに入れる機械語の種類を決める。

cubin は、一個の機械語しか入れられないが、このELFとPTXをくっつけた(多分) fatbin には、 複数のアーキテクチャ用の機械語 + PTXを入れられる。

 $ nvcc -fatbin -arch compute_10 -code=compute_10,sm_20,sm_30 a.cu
 $ cuobjdump a.fatbin
 $ cuobjdump.exe -ptx  -sass a.fatbin

Fatbin elf code:
================
arch = sm_30
code version = [1,4]
producer = cuda
host = windows
compile_size = 32bit
identifier = a.cu

	code for sm_30
		Function : _Z4funcPi
	/*0008*/     /*0x10005de428004001*/ 	MOV R1, c [0x0] [0x44];
	/*0010*/     /*0x00009de428004005*/ 	MOV R2, c [0x0] [0x140];
	/*0018*/     /*0x00201c8580000000*/ 	LD R0, [R2];
	/*0020*/     /*0x04001c034800c000*/ 	IADD R0, R0, 0x1;
	/*0028*/     /*0x00201c8590000000*/ 	ST [R2], R0;
	/*0030*/     /*0x00001de780000000*/ 	EXIT;
	/*0038*/     /*0xe0001de74003ffff*/ 	BRA 0x38;
		..........................



Fatbin elf code:
================
arch = sm_20
code version = [1,4]
producer = cuda
host = windows
compile_size = 32bit
identifier = a.cu

	code for sm_20
		Function : _Z4funcPi
	/*0000*/     /*0x00005de428004404*/ 	MOV R1, c [0x1] [0x100];
	/*0008*/     /*0x80009de428004000*/ 	MOV R2, c [0x0] [0x20];
	/*0010*/     /*0x00201c8580000000*/ 	LD R0, [R2];
	/*0018*/     /*0x04001c034800c000*/ 	IADD R0, R0, 0x1;
	/*0020*/     /*0x00201c8590000000*/ 	ST [R2], R0;
	/*0028*/     /*0x00001de780000000*/ 	EXIT;
		..........................



Fatbin ptx code:
================
arch = sm_10
code version = [1,4]
producer = cuda
host = windows
compile_size = 32bit
identifier = a.cu
ptxasOptions = 
	.version 1.4
.target sm_10, map_f64_to_f32

 // 長いので略
.entry _Z4funcPi (
.param .u32 __cudaparm__Z4funcPi_p)
{
.reg .u32 %r<5>;
.loc	15	2	0
$LDWbegin__Z4funcPi:
.loc	15	3	0
ld.param.u32 %r1, [__cudaparm__Z4funcPi_p];
ld.global.s32 %r2, [%r1+0];
add.s32 %r3, %r2, 1;
st.global.s32 [%r1+0], %r3;
.loc	15	4	0
exit;
$LDWend__Z4funcPi:
} 
 

CUDA コンパイラは、fatbinをランタイムに渡すホストコードを出す。

CUDA ランタイムは、

1. fatbinの中に対応するcubinがあるかを探す。あればそれを実行する
2. fatbinの中にptxがあれば、それを実行時にコンパイルしてそれを実行する

という挙動をする。使いかたとしては、

• 実行時にPTXをアセンブルする時間がもったいない場合は、-code=sm_xxを指定し、対応する機械語を入れておく
• どのマシンで実行するかわからない場合は、-code=compute_xxを指定し、PTXを入れておいて、実行時に機械語を出すようにする

という感じになる。

-arch=sm_xx は、 -arch=compute_xx -code=sm_13,compute_13 と対応している。(なので昔のGPU用にコンパイルしてあCUDAコードは今のGPUでも動く)

アトミック演算や倍精度等の新しい機能を昔のGPUで使う場合は、 -arch=compute_30 -code=sm_10 とか、将来、できるといいなぁ…と、マニュアルには書いてある(今はできない)

まとめ

まとめると、実行時によろしくやってくれるので、-arch sm_10 で特に問題無い。

SASS 命令セット

cuobjdumpのマニュアルに、一応命令一覧と、ひとこと説明が書いてある。

…が、詳しい説明が無くて、どうやったら使えるかわからない命令が多い。

Load and Lock とかあるが、 ll/sc できるのだろうか…

SASS活用事例

ここまで読んだ方は、「これって何かの役に立つの…？」という感が拭えないと思うので、 役に立った事例など紹介。

clock で命令スループット取りたいがptxasが勝手に命令消しやがってこのクソが

命令単位のパフォーマンス取りたい場合は結構あって、NVGPU には、clock() というのがあるし、それが使えるのだが、 ptx は dead code elimination (使ってない命令を消す) 的な最適化をするので、ちゃんと依存関係作ってあげないと測定できない場合が多い。

 $ cat clock.cu
__global__ void f(int n)
{
    __shared__ int p;
    for (int i=0; i<n; i++) {
        int v = 0;
        asm("add.s32 %0, %0, %0;\n\t"
	    "add.s32 %0, %0, %0;\n\t"
	    "add.s32 %0, %0, %0;\n\t"
	    "add.s32 %0, %0, %0;\n\t"
	    :"+r"(v));
        p = v;
    }
}
 $ nvcc -O2 -ptx clock.cu
 $ cat clock.ptx
 // (略)
$Lt_0_1794:
 //<loop> Loop body line 1, nesting depth: 1, estimated iterations: unknown
	.loc	15	6	0
	mov.s32 	%r6, 0;
	mov.u32 	%r7, %r6;
	add.s32 %r7, %r7, %r7;
	add.s32 %r7, %r7, %r7;
	add.s32 %r7, %r7, %r7;
	add.s32 %r7, %r7, %r7;
	
	mov.s32 	%r8, %r7;
	add.s32 	%r4, %r4, 1;
	.loc	15	1	0
	ld.param.s32 	%r1, [__cudaparm__Z1fi_n];
	.loc	15	6	0
	setp.ne.s32 	%p2, %r1, %r4;
	@%p2 bra 	$Lt_0_1794;
 // (略)

PTXでは出てるのに、実際測ってみると何もとれない

SASS で見ると消えてる…

$ cuobjdump.exe -sass clock.cubin

	code for sm_10
		Function : _Z1fi
	/*0000*/     /*0x307cc9fd6c20c7c8*/ 	ISET.S32.C0 o [0x7f], g [0x4], R124, LE;
	/*0008*/     /*0x3000000300000280*/ 	RET C0.NE;
	/*0010*/     /*0x1000f8010403c780*/ 	MOV R0, R124;
	/*0018*/     /*0x2001800100000003*/ 	IADD32I R0, R0, 0x1;
	/*0020*/     /*0x3000c9fd6c2147c8*/ 	ISET.S32.C0 o [0x7f], g [0x4], R0, NE;
	/*0028*/     /*0x1000300300000280*/ 	BRA C0.NE, 0x18;
	/*0030*/     /*0x00000a01e43f0781*/ 	R2G.U32.U32 g [0x5], R124;
		......................

入力をメモリにするとちゃんと出る

 $ cat clock.cu
__global__ void f(int n)
{
    __shared__ int p;
    for (int i=0; i<n; i++) {
        int v = p;
        asm("add.s32 %0, %0, %0;\n\t"
	    "add.s32 %0, %0, %0;\n\t"
	    "add.s32 %0, %0, %0;\n\t"
	    "add.s32 %0, %0, %0;\n\t"
	    :"+r"(v));
        p = v;
    }
}

 $ cuobjdump.exe -sass clock.cubin

	code for sm_10
		Function : _Z1fi
	/*0000*/     /*0x307cc9fd6c20c7c8*/ 	ISET.S32.C0 o [0x7f], g [0x4], R124, LE;
	/*0008*/     /*0x3000000300000280*/ 	RET C0.NE;
	/*0010*/     /*0x1000ca010423c780*/ 	MOV R0, g [0x5];
	/*0018*/     /*0x1000f8050403c780*/ 	MOV R1, R124;
	/*0020*/     /*0x20008000        */ 	IADD32 R0, R0, R0;
	/*0024*/     /*0x20008000        */ 	IADD32 R0, R0, R0;
	/*0028*/     /*0x2001820500000003*/ 	IADD32I R1, R1, 0x1;
	/*0030*/     /*0x2000000104000780*/ 	IADD R0, R0, R0;
	/*0038*/     /*0x3001c9fd6c2147c8*/ 	ISET.S32.C0 o [0x7f], g [0x4], R1, NE;
	/*0040*/     /*0x2000000104000780*/ 	IADD R0, R0, R0;
	/*0048*/     /*0x1000400300000280*/ 	BRA C0.NE, 0x20;
	/*0050*/     /*0x00000a01e4200781*/ 	R2G.U32.U32 g [0x5], R0;
		......................

こういうのを確認するときにcuobjdump必須

add.u64 (GT200)

64bitマシンでは、アドレスの足し算が64bit加算になるが、NVGPUは 64bit加算が無くて2命令になってしまう。

__global__ void f(int *p, int *q, int off, int n)
{
    for (int i=0; i<n; i++) {
	int *ptr = &p[off+i];
	if (ptr < q) {
	    *ptr += 1;
	}
    }
}
↓
	/*0020*/     /*0x2000000504004780*/ 	IADD R1, R0, R1;
	/*0028*/     /*0x2000c809042007c0*/ 	IADD.C0 R2, g [0x4], R0;
	/*0030*/     /*0x3040ca0d043f0780*/ 	IADD.CARRY0 R3, g [0x5], R124;

が、GT200 は、GPU側のアドレス空間が32bitしか無いので、実際に64bit加算が必要無い場合が多い。 ptxas は涙ぐましい努力をして、ptx の add.u64 を 32bit 加算にする努力をしているように見える。

__global__ void f2(int *p, int *q, int off, int n)
{
    for (int i=0; i<n; i++) {
	int *ptr = &p[off+i];
	*ptr += 1;
    }
}

	/*0020*/     /*0x2000000504004780*/ 	IADD R1, R0, R1;
	/*0028*/     /*0x2000c80d04200780*/ 	IADD R3, g [0x4], R0;
	/*0030*/     /*0xd00e060980c00780*/ 	GLD.U32 R2, global14 [R3];

こういうのを確認するのに使える。

ただ残念ながらsm_20以降はGPUアドレス空間も64bitになって、涙ぐましい努力は行われなくなっている

	/*00c8*/     /*0x04009c034800c000*/ 	IADD R2, R0, 0x1;
	/*00d0*/     /*0x20411c034801c000*/ 	IADD R4.CC, R4, 0x8;
	/*00d8*/     /*0x5021dc2318804005*/ 	ISETP.LT.AND P0, pt, R2, c [0x0] [0x154], P0;
	/*00e0*/     /*0xfc515c4348000000*/ 	IADD.X R5, R5, RZ;
	/*00e8*/     /*0x200001e74003fffe*/ 	@P0 BRA 0x78;
	/*00f0*/     /*0x5001dc23188e4005*/ 	ISETP.LT.AND P0, pt, R0, c [0x0] [0x154], pt;
	/*00f8*/     /*0x000021e780000000*/ 	@!P0 EXIT;
	/*0108*/     /*0x00401c8584000000*/ 	LD.E R0, [R4];

メモリモデルと __threadfence_block() の効果

TeslaC2050(Fermi世代)がでたころ、社内でsm_20だとsumがちゃんと動かないんだけど…とかメールがあった。

原因としては、reductionするときのメモリバリアをちゃんと書いていなかったというものだった。

当時のCUDAマニュアルには、"__threadfence_block()書かないとコンパイラがメモリアクセス消しちゃうからちゃんと__threadfence_block()書いてね" とか書いてあるが、その現象では、コンパイラはちゃんとコード出してた、が、__threadfence_block()を入れると問題が解決した。

→ __threadfence_block() は、コンパイラの最適化を抑制する以上の効果があるのか？

というときも、SASSを見れば一発

// __global__ void f() { __threadfence_block(); }

Fatbin elf code:
================
arch = sm_20
code version = [1,4]
producer = cuda
host = windows
compile_size = 32bit
identifier = tf.cu

	code for sm_20
		Function : _Z2tfv
	/*0000*/     /*0x00005de428004404*/ 	MOV R1, c [0x1] [0x100];
	/*0008*/     /*0x00001c05e0000000*/ 	MEMBAR.CTA;
	/*0010*/     /*0x00000007d0000000*/ 	BPT.DRAIN 0x0;
	/*0018*/     /*0x00001de780000000*/ 	EXIT;
		.......................



Fatbin elf code:
================
arch = sm_10
code version = [1,2]
producer = cuda
host = windows
compile_size = 32bit
identifier = tf.cu

	code for sm_10
		Function : _Z2tfv
	/*0000*/     /*0xf0000001e0000001*/ 	NOP;
		.......................

sm_10 では、何もしていないが、sm_10 では、member.cta というのが出ている。→ sm_20では、なんかメモリモデル変だからメモリバリア必要か、と推測できる

レジスタ数

(自分ではここまで追いつめられたことないのでよく知らない)

CUDAではレジスタ数、スレッド数のバランスが結構センシティブなのだが、 ptxas の maxreg で調整するとスケジューリング悪くなって、でもPTXだとレジスタよくわからなくて…とかになる(らしい)

SASSは実レジスタが見えるので、そういう時もどこでレジスタ足りなくなってるか、とか確認できる

2ch の トリップ検索スレによると実際に確認してる人がいるらしい

492 ： ◆MERIKEN4.k ：2011/09/11(日) 01:06:24.80 ID:tGzUJSSY0
    >>490
    SASS見てみましたけど、予想通りS-Boxのあたりでレジスタを大量に消費してました。
    30以上使ってたので、実際に必要な数よりかなり多いですね。
    cubinとPTXのコードが違うというのはもちろん理解してますけど、
    実際の挙動はPTXのコードとプロファイリングの結果からある程度予測できるという印象です。

    >>491
    自分の調べた限りでは、現在のBitslice DESの最小ゲート数の実装は、John the Ripperで
    使われているRoman Rusakov氏のものでした。>>323のSlashdotの記事で紹介されてますけど、
    Kahn氏のオリジナルより17%ゲート数が少ないとのことです。 

まとめ

PTXよりも情報量あるので、PTX読むくらいならSASS読みましょう

この文章について

この文書は関東GPGPU勉強会のために書かれました。