ptx/src/test/spirv_run/call.ptx


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

.version 6.5
.target sm_30
.address_size 64

.visible .func (.param.u64 output) incr (.param.u64 input);

.visible .entry call(
	.param .u64 input,
	.param .u64 output
)
{
	.reg .u64 	    in_addr;
    .reg .u64 	    out_addr;
    .reg .u64 	    temp;

	ld.param.u64 	in_addr, [input];
    ld.param.u64 	out_addr, [output];

    ld.global.u64   temp, [in_addr];
    .param.u64      incr_in;
    .param.u64      incr_out;
    st.param.b64    [incr_in], temp;
    call (incr_out), incr, (incr_in);
    ld.param.u64    temp, [incr_out];
    st.global.u64   [out_addr], temp;
    ret;
}

.visible .func (.param .u64 output) incr(
    .param .u64 input
)
{
    .reg .u64 	    temp;
    ld.param.u64    temp, [input];
    add.u64         temp, temp, 1;
    st.param.u64    [output], temp;
    ret;
}