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

.version 6.5
.target sm_30
.address_size 64

.extern .shared .b32 shared_ex[];
.shared .b32 shared_mod[4];


.func (.reg .b64 out) load_from_shared()
{
    ld.shared.u64   out, [shared_mod];
    ret;
}

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

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

    ld.u64          temp1, [in_addr];
    st.shared.u64   [shared_ex], temp1;
    call            (temp2), load_from_shared;
	add.u64		    temp2, temp2, temp1;
    st.u64          [out_addr], temp2;
	ret;
}