rtdevlib: fix function signature mismatches

A issue discovered while running HIP programs on OpenCL-BE->rusticl
(with #830). Linking of __chip_atomic_add_f* symbols failed because
the caller's and callee's function signature differed by their pointer
parameters (pointee type didn't match).

The mismatch was caused by LLVM-SPIRV-Translator's feature that
attempts to recover original pointee types in LLVM bitcodes that use
opaque pointers. But the way it attempts to infer the types may end up
with SPIR-V functions with different pointee type across SPIR-V
modules.

The issue is worked around by passing pointers as integers for
functions whose definitions are linked in at runtime.

Author: Henry Linjamäki

bitcode/README-devicelib.md

@@ -35,4 +35,23 @@ Note: Some amdgcn intrinsics still don't have generic equivalents so for example
 * If there are missing headers, add them to the end of the list with a comment that they're undocumented. The missing documentationt should be reported to AMD or NVIDIA. 
 * Avoid using macros, escpecially in the headers.
 * `devicelib.cl` does use some macros mainly to avoid code duplication when declaring function overloads. If you do use macros, make sure to leave a comment with the full function name that the macro expands to.
-* All device-side functions should be declared with `extern "C"` and follow the following naming convention: `__chip_<function_name_snake_case>_<type>` where type is one of `f32`, `f64`, `i32`, `i64`, `u32`, `u64`. For example, `__chip_sin_f32` is the device-side function that implements the `sin` function for `float` arguments. Note: this convention is not yet fully implemented.
+* All device-side functions should be declared with `extern "C"` and follow the following naming convention: `__chip_<function_name_snake_case>_<type>` where type is one of `f32`, `f64`, `i32`, `i64`, `u32`, `u64`. For example, `__chip_sin_f32` is the device-side function that implements the `sin` function for `float` arguments. Note: this convention is not yet fully implemented.
+
+## Obfuscated Pointer Parameters
+
+A selection of devicelib definitions use `__chip_obfuscated_ptr_t` as
+parameter type insted of regular pointer type. This is for working
+around LLVM-SPIR-VTranslation feature that attempts to recover
+original pointee types in LLVM bitcode input which uses opaque
+pointers (only option with the latest LLVM).
+
+An issue with the feature is that the type inference is influenced by
+the surrounding code and this causes situations where same LLVM
+function declarations and definitions may end up to be have parameters
+with different pointee types across SPIR-V modules. Linking such
+modules together will probably trigger undefined behavior. Al least on
+Mesa's rusticl frontend the linking is known to fail.
+
+The shortcoming of the type inference is worked around by passing the
+pointer arguments as some non-pointer type (`__chip_obfuscated_ptr_t`)
+and it's used for devicelib definitions which are linked at runtime.

bitcode/atomicAddDouble_emulation.cl

@@ -22,6 +22,8 @@
 
 // Implementations for emulated 64-bit floating point atomic add operations.
 
+#include "cl_utils.h"
+
 #ifndef __opencl_c_generic_address_space
 #error __opencl_c_generic_address_space needed!
 #endif
@@ -57,16 +59,17 @@ static OVERLOADED double __chip_atomic_add_f64(volatile global double *address,
   return as_double(r);
 }
 
-double __chip_atomic_add_f64(generic double *address, double val) {
-  volatile global double *gi = to_global(address);
+double __chip_atomic_add_f64(__chip_obfuscated_ptr_t address, double val) {
+  volatile global double *gi = to_global(UNCOVER_OBFUSCATED_PTR(address));
   if (gi)
     return __chip_atomic_add_f64(gi, val);
-  volatile local double *li = to_local(address);
+  volatile local double *li = to_local(UNCOVER_OBFUSCATED_PTR(address));
   if (li)
     return __chip_atomic_add_f64(li, val);
   return 0;
 }
 
-double __chip_atomic_add_system_f64(generic double *address, double val) {
+double __chip_atomic_add_system_f64(__chip_obfuscated_ptr_t address,
+                                    double val) {
   return __chip_atomic_add_f64(address, val);
 }

bitcode/atomicAddDouble_native.cl

@@ -23,6 +23,8 @@
 // Implementations for 64-bit floating point atomic operations using
 // OpenCL built-in extension.
 
+#include "cl_utils.h"
+
 #ifndef __opencl_c_generic_address_space
 #error __opencl_c_generic_address_space needed!
 #endif
@@ -37,9 +39,11 @@
 /* https://registry.khronos.org/OpenCL/extensions/ext/cl_ext_float_atomics.html
  */
 #define DEF_CHIP_ATOMIC2F_ORDER_SCOPE(NAME, OP, ORDER, SCOPE)                  \
-  double __chip_atomic_##NAME##_f64(double *address, double i) {               \
-    return atomic_##OP##_explicit((volatile __generic double *)address, i,     \
-                                  memory_order_##ORDER, memory_scope_##SCOPE); \
+  double __chip_atomic_##NAME##_f64(__chip_obfuscated_ptr_t address,           \
+                                    double i) {                                \
+    return atomic_##OP##_explicit(                                             \
+        (volatile __generic double *)UNCOVER_OBFUSCATED_PTR(address), i,       \
+        memory_order_##ORDER, memory_scope_##SCOPE);                           \
   }
 
 #define DEF_CHIP_ATOMIC2F(NAME, OP)                                            \

bitcode/atomicAddFloat_emulation.cl

@@ -22,6 +22,8 @@
 
 // Implementations for emulated 32-bit floating point atomic add operations.
 
+#include "cl_utils.h"
+
 #ifndef __opencl_c_generic_address_space
 #error __opencl_c_generic_address_space needed!
 #endif
@@ -57,16 +59,16 @@ static OVERLOADED float __chip_atomic_add_f32(volatile global float *address,
   return as_float(r);
 }
 
-float __chip_atomic_add_f32(generic float *address, float val) {
-  volatile global float *gi = to_global(address);
+float __chip_atomic_add_f32(__chip_obfuscated_ptr_t address, float val) {
+  volatile global float *gi = to_global(UNCOVER_OBFUSCATED_PTR(address));
   if (gi)
     return __chip_atomic_add_f32(gi, val);
-  volatile local float *li = to_local(address);
+  volatile local float *li = to_local(UNCOVER_OBFUSCATED_PTR(address));
   if (li)
     return __chip_atomic_add_f32(li, val);
   return 0;
 }
 
-float __chip_atomic_add_system_f32(generic float *address, float val) {
+float __chip_atomic_add_system_f32(__chip_obfuscated_ptr_t address, float val) {
   return __chip_atomic_add_f32(address, val);
 }

bitcode/atomicAddFloat_native.cl

@@ -23,6 +23,8 @@
 // Implementations for 32-bit floating point atomic operations using
 // OpenCL built-in extension.
 
+#include "cl_utils.h"
+
 #ifndef __opencl_c_generic_address_space
 #error __opencl_c_generic_address_space needed!
 #endif
@@ -37,9 +39,10 @@
 /* https://registry.khronos.org/OpenCL/extensions/ext/cl_ext_float_atomics.html
  */
 #define DEF_CHIP_ATOMIC2F_ORDER_SCOPE(NAME, OP, ORDER, SCOPE)                  \
-  float __chip_atomic_##NAME##_f32(float *address, float i) {                  \
-    return atomic_##OP##_explicit((volatile __generic float *)address, i,      \
-                                  memory_order_##ORDER, memory_scope_##SCOPE); \
+  float __chip_atomic_##NAME##_f32(__chip_obfuscated_ptr_t address, float i) { \
+    return atomic_##OP##_explicit(                                             \
+        (volatile __generic float *)UNCOVER_OBFUSCATED_PTR(address), i, \
+        memory_order_##ORDER, memory_scope_##SCOPE);                           \
   }
 
 #define DEF_CHIP_ATOMIC2F(NAME, OP)                                            \

bitcode/cl_utils.h

@@ -0,0 +1,32 @@
+/*
+ * Copyright (c) 2024 chipStar developers
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
+ * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
+ * DEALINGS IN THE SOFTWARE.
+ */
+
+// Shared OpenCL utilities.
+
+#ifndef BITCODE_CL_UTILS_H
+#define BITCODE_CL_UTILS_H
+
+// See bitcode/README-devicelib.md for the purpose of "obfuscated pointers".
+typedef ulong __chip_obfuscated_ptr_t;
+#define UNCOVER_OBFUSCATED_PTR(_PTR) ((generic void *)_PTR)
+
+#endif // BITCODE_CL_UTILS_H

include/hip/devicelib/atomics.hh

@@ -26,6 +26,13 @@
 #include <hip/host_defines.h>
 #include <hip/devicelib/macros.hh>
 
+// See bitcode/README-devicelib.md for the purpose of "obfuscated pointers".
+typedef unsigned long __chip_obfuscated_ptr_t;
+inline __device__ __chip_obfuscated_ptr_t __chip_obfuscate_ptr(void *ptr) {
+  static_assert(sizeof(__chip_obfuscated_ptr_t) == sizeof(void *), "");
+  return reinterpret_cast<__chip_obfuscated_ptr_t>(ptr);
+}
+
 // Copied from HIP programming guide:
 // https://docs.amd.com/bundle/HIP-Programming-Guide-v5.0/page/Programming_with_HIP.html
 // Slightly modified to group operations
@@ -63,14 +70,16 @@ atomicAdd(unsigned long long *address, unsigned long long val) {
   return __chip_atomic_add_l(address, val);
 }
 
-extern "C" __device__ float __chip_atomic_add_f32(float *address, float val);
+extern "C" __device__ float
+__chip_atomic_add_f32(__chip_obfuscated_ptr_t address, float val);
 extern "C++" inline __device__ float atomicAdd(float *address, float val) {
-  return __chip_atomic_add_f32(address, val);
+  return __chip_atomic_add_f32(__chip_obfuscate_ptr(address), val);
 }
 
-extern "C" __device__ double __chip_atomic_add_f64(double *address, double val);
+extern "C" __device__ double
+__chip_atomic_add_f64(__chip_obfuscated_ptr_t address, double val);
 extern "C++" inline __device__ double atomicAdd(double *address, double val) {
-  return __chip_atomic_add_f64(address, val);
+  return __chip_atomic_add_f64(__chip_obfuscate_ptr(address), val);
 }
 
 extern "C" __device__ int __chip_atomic_add_system_i(int *address, int val);
@@ -92,18 +101,18 @@ atomicAdd_system(unsigned long long *address, unsigned long long val) {
   return __chip_atomic_add_system_l(address, val);
 }
 
-extern "C" __device__ float __chip_atomic_add_system_f32(float *address,
-                                                         float val);
+extern "C" __device__ float
+__chip_atomic_add_system_f32(__chip_obfuscated_ptr_t address, float val);
 extern "C++" inline __device__ float atomicAdd_system(float *address,
                                                       float val) {
-  return __chip_atomic_add_system_f32(address, val);
+  return __chip_atomic_add_system_f32(__chip_obfuscate_ptr(address), val);
 }
 
-extern "C" __device__ double __chip_atomic_add_system_f64(double *address,
-                                                          double val);
+extern "C" __device__ double
+__chip_atomic_add_system_f64(__chip_obfuscated_ptr_t address, double val);
 extern "C++" inline __device__ double atomicAdd_system(double *address,
                                                        double val) {
-  return __chip_atomic_add_system_f64(address, val);
+  return __chip_atomic_add_system_f64(__chip_obfuscate_ptr(address), val);
 }
 
 extern "C" __device__ int __chip_atomic_sub_i(int *address, int val);
@@ -432,7 +441,7 @@ atomicXor_system(unsigned long long *address, unsigned long long val) {
 
 // Undocumented
 extern "C++" inline __device__ void atomicAddNoRet(float *address, float val) {
-  (void)__chip_atomic_add_f32(address, val);
+  (void)__chip_atomic_add_f32(__chip_obfuscate_ptr(address), val);
 }
 
 #endif // HIP_INLUDE_DEVICELIB_ATOMICS