example of hacking spirv64-amd into this

Author: JonChesterfield

clang/lib/Headers/gpuintrin.h

@@ -56,14 +56,21 @@ __gpu_match_all_u64_impl(uint64_t __lane_mask, uint64_t __x);
 _Pragma("omp end declare variant");
 _Pragma("omp end declare target");
 
+
+
 #if defined(__NVPTX__)
-#include <nvptxintrin.h>
+  #include <nvptxintrin.h>
 #elif defined(__AMDGPU__)
-#include <amdgpuintrin.h>
+  #if defined(__SPIRV64__)
+    // the spirv64-amd-amdhsa triple claims to be amdgpu and spirv
+    #include <spirvintrin.h>
+  #else
+    #include <amdgpuintrin.h>
+  #endif
 #elif defined(__SPIRV64__)
-#include <spirvintrin.h>
+  #include <spirvintrin.h>
 #elif !defined(_OPENMP)
-#error "This header is only meant to be used on GPU architectures."
+  #error "This header is only meant to be used on GPU architectures."
 #endif
 
 _Pragma("omp begin declare target device_type(nohost)");

clang/lib/Headers/spirvintrin.h

@@ -37,7 +37,13 @@
 #define __gpu_kernel
 
 // Returns the number of workgroups in the 'x' dimension of the grid.
-_DEFAULT_FN_ATTRS uint32_t __gpu_num_blocks_x(void);
+_DEFAULT_FN_ATTRS uint32_t __gpu_num_blocks_x(void) {
+#if defined(__AMDGPU__)
+  return __builtin_amdgcn_grid_size_x() / __builtin_amdgcn_workgroup_size_x();
+#else
+#error "Not yet implemented"
+#endif
+}
 
 // Returns the number of workgroups in the 'y' dimension of the grid.
 _DEFAULT_FN_ATTRS uint32_t __gpu_num_blocks_y(void);
@@ -83,7 +89,12 @@ _DEFAULT_FN_ATTRS uint64_t __gpu_lane_mask(void);
 
 // Copies the value from the first active thread in the wave to the rest.
 _DEFAULT_FN_ATTRS uint32_t __gpu_read_first_lane_u32(uint64_t __lane_mask,
-                                                     uint32_t __x);
+                                                     uint32_t __x) {
+
+  uint64_t __gpu_read_first_lane_u32_impl(
+      uint64_t __lane_mask, uint32_t __x) asm("llvm.spv.wave.readlane.i64");
+  return __gpu_read_first_lane_u32_impl(__lane_mask, __x);
+}
 
 // Returns a bitmask of threads in the current lane for which \p x is true.
 _DEFAULT_FN_ATTRS uint64_t __gpu_ballot(uint64_t __lane_mask, bool __x);