AMDGCN inefficient long add with constant

[preda]

Consider this OpenCL kernel:

 kernel void testAdd(global long* io) {
  long C = ((long) 1) << 50;
  io[get_global_id(0)] = C + io[get_global_id(0)];
}

This ISA is generated for the long add:

	v_mov_b32_e32 v4, 0x40000
	v_add_co_u32_e32 v2, vcc, 0, v2
	v_addc_co_u32_e32 v3, vcc, v3, v4, vcc

As you see, the above code is.. un-necessary. In particular, v_add_co_u32_e32 v2, vcc, 0, v2 does not change the value of v2, and can not produce carry-out.

[preda]

Expected would be something like:

v_mov_b32_e32 v4, 0x40000
v_add_co_u32_e32 v3, vcc, v3, v4

[preda]

And a small observation: the same code is generated and the problem is easier to see if "long" is replaced with "unsigned long" in the sample kernel.

[ppanchad-amd]

Hi @preda. Internal ticket has been created to investigate your issue. Thanks!

[sohaibnd]

Hi @preda, what options did you use to compile that code?

[preda]

Hi @preda, what options did you use to compile that code?

@sohaibnd

-cl-finite-math-only -cl-std=CL2.0

[sohaibnd]

@preda This is an optimization bug, thanks for pointing it out. See llvm#122049 for the fix being put in.