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

// HACK ALERT
// This test is for testing workaround for a bug in IGC where linking fails
// if there is shl/shr with different width of value and shift

.version 6.5
.target sm_30
.address_size 64

.visible .entry shl_link_hack(
	.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];

    // Here only to trigger linking
    .reg .u32 	    unused;
    atom.inc.u32    unused, [out_addr], 2000000;

    ld.u64          temp, [in_addr];
	shl.b64		    temp2, temp, 2;
    st.u64          [out_addr], temp2;
	ret;
}