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

.version 6.5
.target sm_30
.address_size 64

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

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

	atom.inc.u32	    temp1, [in_addr], 101;
    atom.global.inc.u32 temp2, [in_addr], 101;
    ld.u32              temp3, [in_addr];
    st.u32              [out_addr], temp1;
    st.u32              [out_addr+4], temp2;
    st.u32              [out_addr+8], temp3;
	ret;
}