OpenCL behavior when re-defining builtin math functions

[preda]

Consider the following example, which implements a complex multiply-add function named "mad" with a signature that clashes with the builtin OpenCL math function mad(gentype, gentype, gentype) (where gentype includes double2).

__attribute__((overloadable)) double2 mad(double2 a, double2 b, double2 c) {
  return (double2) (fma(a.x, b.x, fma(a.y, -b.y, c.x)), fma(a.x, b.y, fma(a.y, b.x, c.y)));
}

kernel void testKernel(global double2* out, global double2* in) {
  uint me = get_local_id(0);

  double2 a = in[3*me];
  double2 b = in[3*me + 1];
  double2 c = in[3*me + 2];
  out[me] = mad(a, b, c);
}

This example is compiled very differently by ROCm 6.1.3 and ROCm 6.2.0. In both case it is accepted without errors or warnings. The resulting compilation is: With ROCm 6.1.3, the user-defined mad() is used, with ROCm 6.2.0, the built-in mad() is used.

In particular the 6.2.0 behavior, where the user defined function is accepted without errors and silently ignored afterwards, is not acceptable IMO and represents a bug.

[schung-amd]

Hi @preda, sorry for the delay on this, and sorry it broke your code. I've taken a look into what causes this, my best guess is that this is similar to ELF interposition. The linker is implicitly giving the built-ins precedence due to linking order (the specifics here are not important, we treat this as UB so there are no guarantees of consistency, and the compiler doesn't care as long as the symbol is defined somewhere. No errors are thrown by the linker because libraries are allowed to define symbols identically to other libraries. The difficulty in this case is that the built-ins are being linked in under the hood during runtime compilation and the user has no control over the linking order.

The solution is, as you've found, to not define functions with identical names and signatures to built-ins. Of course, in your case this was not an issue originally and there was no indication that this would be an issue, and I apologize for the inconvenience. I'd also like to note that while the built-ins are defined with __attribute__((overloadable)), this is really for our own purposes to overload the built-ins for various types, and was not meant to enable users to overload the built-ins. From our perspective defining a function with the same signature as a built-in is UB.