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

.version 6.5
.target sm_30
.address_size 64

.extern .shared .align 4 .b8 shared_mem[];

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

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

    ld.global.u64           temp1, [in_addr];
    st.shared.u64           [shared_addr], temp1;
    ld.shared.u64           temp2, [shared_addr];
    st.global.u64           [out_addr], temp2;
	ret;
}