ptx/src/test/spirv_run/extern_func.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
39

.version 6.5
.target sm_30
.address_size 64

.extern .func  (.param .align 16 .b8 func_retval0[16]) foobar
(
	.param .b64 extern_func
)
;

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

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

    ld.global.u64   temp, [in_addr];

	{
        .param .b64 param0;
        st.param.b64	[param0+0], temp;
        .param .align 16 .b8 retval0[16];
        call.uni (retval0) ,
        foobar ,
        (
        param0
        );
        ld.param.u64    temp2, [retval0];
	}
    st.u64          [out_addr], temp2;
	ret;
}