decuda/gencvt.py at master · corinus/decuda · GitHub

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
#!/usr/bin/python
from AutoComp import PTX,CompilationError
import sys

ptx = """
.version 1.0
.target compute_10, map_f64_to_f32

.entry bra_test
{
.reg .%(src)s $r1;
.reg .u64 $rd1,$rd2;
.reg .%(dst)s $f1;
.param .u64 __cudaparm_data;
.param .u64 __cudaparm_data2;

ld.param.u64 	$rd1, [__cudaparm_data2];
ld.global.%(src)s 	$r1, [$rd1+0];
cvt%(round)s.%(dst)s.%(src)s 	$f1, $r1;
ld.param.u64 	$rd2, [__cudaparm_data];
st.global.%(dst)s 	[$rd2+0], $f1;

exit;
}
"""


types = ["u8","u16","u32","s8","s16","s32","f32"] #"f16", "f64","u64","s64"
modes_fp = [".rn", ".rz", ".rm", ".rp"]
modes_int = [".rni", ".rzi", ".rmi", ".rpi"]

#vars = {"src":"s32", "dst":"f32", "round":".rn"}
#x = PTX(ptx % vars)
#disa = x.bin.kernels[0].disassemble()

for src in types:
    for dst in types:
        modes = [""]
        if dst[0]=="f":
            modes = modes_fp
        else:
            modes = modes_int
        for mode in modes:
            print src,dst,mode
            vars = {"src":src, "dst":dst, "round":mode}
            print src, dst, mode
            try:
                x = PTX(ptx % vars)
                disa = x.bin.kernels[0].disassemble(sys.stdout)
                #lines = disa.split("\n")
                #if len(lines)>5:
                #    print lines[3]
            except CompilationError,e:
                print e.message
            print "====================="