c89 support

Author: jhuber6

clang/lib/Headers/amdgpuintrin.h

@@ -13,15 +13,11 @@
 #error "This file is intended for AMDGPU targets or offloading to AMDGPU"
 #endif
 
-#include <stdbool.h>
 #include <stdint.h>
 
-#if !defined(_DEFAULT_FN_ATTRS)
-#if defined(__HIP__) || defined(__CUDA__)
-#define _DEFAULT_FN_ATTRS __attribute__((device))
-#else
-#define _DEFAULT_FN_ATTRS
-#endif
+#if !defined(__cplusplus)
+#pragma push_macro("bool")
+#define bool _Bool
 #endif
 
 #pragma omp begin declare target device_type(nohost)
@@ -38,118 +34,122 @@
 #define __gpu_kernel __attribute__((amdgpu_kernel, visibility("protected")))
 
 // Returns the number of workgroups in the 'x' dimension of the grid.
-_DEFAULT_FN_ATTRS static inline uint32_t __gpu_num_blocks_x(void) {
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_blocks_x(void) {
   return __builtin_amdgcn_grid_size_x() / __builtin_amdgcn_workgroup_size_x();
 }
 
 // Returns the number of workgroups in the 'y' dimension of the grid.
-_DEFAULT_FN_ATTRS static inline uint32_t __gpu_num_blocks_y(void) {
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_blocks_y(void) {
   return __builtin_amdgcn_grid_size_y() / __builtin_amdgcn_workgroup_size_y();
 }
 
 // Returns the number of workgroups in the 'z' dimension of the grid.
-_DEFAULT_FN_ATTRS static inline uint32_t __gpu_num_blocks_z(void) {
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_blocks_z(void) {
   return __builtin_amdgcn_grid_size_z() / __builtin_amdgcn_workgroup_size_z();
 }
 
 // Returns the 'x' dimension of the current AMD workgroup's id.
-_DEFAULT_FN_ATTRS static inline uint32_t __gpu_block_id_x(void) {
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_block_id_x(void) {
   return __builtin_amdgcn_workgroup_id_x();
 }
 
 // Returns the 'y' dimension of the current AMD workgroup's id.
-_DEFAULT_FN_ATTRS static inline uint32_t __gpu_block_id_y(void) {
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_block_id_y(void) {
   return __builtin_amdgcn_workgroup_id_y();
 }
 
 // Returns the 'z' dimension of the current AMD workgroup's id.
-_DEFAULT_FN_ATTRS static inline uint32_t __gpu_block_id_z(void) {
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_block_id_z(void) {
   return __builtin_amdgcn_workgroup_id_z();
 }
 
 // Returns the number of workitems in the 'x' dimension.
-_DEFAULT_FN_ATTRS static inline uint32_t __gpu_num_threads_x(void) {
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_threads_x(void) {
   return __builtin_amdgcn_workgroup_size_x();
 }
 
 // Returns the number of workitems in the 'y' dimension.
-_DEFAULT_FN_ATTRS static inline uint32_t __gpu_num_threads_y(void) {
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_threads_y(void) {
   return __builtin_amdgcn_workgroup_size_y();
 }
 
 // Returns the number of workitems in the 'z' dimension.
-_DEFAULT_FN_ATTRS static inline uint32_t __gpu_num_threads_z(void) {
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_threads_z(void) {
   return __builtin_amdgcn_workgroup_size_z();
 }
 
 // Returns the 'x' dimension id of the workitem in the current AMD workgroup.
-_DEFAULT_FN_ATTRS static inline uint32_t __gpu_thread_id_x(void) {
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_thread_id_x(void) {
   return __builtin_amdgcn_workitem_id_x();
 }
 
 // Returns the 'y' dimension id of the workitem in the current AMD workgroup.
-_DEFAULT_FN_ATTRS static inline uint32_t __gpu_thread_id_y(void) {
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_thread_id_y(void) {
   return __builtin_amdgcn_workitem_id_y();
 }
 
 // Returns the 'z' dimension id of the workitem in the current AMD workgroup.
-_DEFAULT_FN_ATTRS static inline uint32_t __gpu_thread_id_z(void) {
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_thread_id_z(void) {
   return __builtin_amdgcn_workitem_id_z();
 }
 
 // Returns the size of an AMD wavefront, either 32 or 64 depending on hardware
 // and compilation options.
-_DEFAULT_FN_ATTRS static inline uint32_t __gpu_num_lanes(void) {
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_lanes(void) {
   return __builtin_amdgcn_wavefrontsize();
 }
 
 // Returns the id of the thread inside of an AMD wavefront executing together.
-_DEFAULT_FN_ATTRS static inline uint32_t __gpu_lane_id(void) {
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_lane_id(void) {
   return __builtin_amdgcn_mbcnt_hi(~0u, __builtin_amdgcn_mbcnt_lo(~0u, 0u));
 }
 
 // Returns the bit-mask of active threads in the current wavefront.
-_DEFAULT_FN_ATTRS static inline uint64_t __gpu_lane_mask(void) {
+_DEFAULT_FN_ATTRS static __inline__ uint64_t __gpu_lane_mask(void) {
   return __builtin_amdgcn_read_exec();
 }
 
 // Copies the value from the first active thread in the wavefront to the rest.
-_DEFAULT_FN_ATTRS static inline uint32_t
+_DEFAULT_FN_ATTRS static __inline__ uint32_t
 __gpu_broadcast_u32(uint64_t __lane_mask, uint32_t __x) {
   return __builtin_amdgcn_readfirstlane(__x);
 }
 
 // Returns a bitmask of threads in the current lane for which \p x is true.
-_DEFAULT_FN_ATTRS static inline uint64_t __gpu_ballot(uint64_t __lane_mask,
-                                                      bool __x) {
+_DEFAULT_FN_ATTRS static __inline__ uint64_t __gpu_ballot(uint64_t __lane_mask,
+                                                          bool __x) {
   // The lane_mask & gives the nvptx semantics when lane_mask is a subset of
   // the active threads
   return __lane_mask & __builtin_amdgcn_ballot_w64(__x);
 }
 
 // Waits for all the threads in the block to converge and issues a fence.
-_DEFAULT_FN_ATTRS static inline void __gpu_sync_threads(void) {
+_DEFAULT_FN_ATTRS static __inline__ void __gpu_sync_threads(void) {
   __builtin_amdgcn_s_barrier();
   __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup");
 }
 
 // Wait for all threads in the wavefront to converge, this is a noop on AMDGPU.
-_DEFAULT_FN_ATTRS static inline void __gpu_sync_lane(uint64_t __lane_mask) {
+_DEFAULT_FN_ATTRS static __inline__ void __gpu_sync_lane(uint64_t __lane_mask) {
   __builtin_amdgcn_wave_barrier();
 }
 
 // Shuffles the the lanes inside the wavefront according to the given index.
-_DEFAULT_FN_ATTRS static inline uint32_t
+_DEFAULT_FN_ATTRS static __inline__ uint32_t
 __gpu_shuffle_idx_u32(uint64_t __lane_mask, uint32_t __idx, uint32_t __x) {
   return __builtin_amdgcn_ds_bpermute(__idx << 2, __x);
 }
 
 // Terminates execution of the associated wavefront.
-_DEFAULT_FN_ATTRS [[noreturn]] static inline void __gpu_exit(void) {
+_DEFAULT_FN_ATTRS [[noreturn]] static __inline__ void __gpu_exit(void) {
   __builtin_amdgcn_endpgm();
 }
 
 #pragma omp end declare variant
 #pragma omp end declare target
 
+#if !defined(__cplusplus)
+#pragma pop_macro("bool")
+#endif
+
 #endif // __AMDGPUINTRIN_H

clang/lib/Headers/gpuintrin.h

@@ -17,6 +17,14 @@
 #ifndef __GPUINTRIN_H
 #define __GPUINTRIN_H
 
+#if !defined(_DEFAULT_FN_ATTRS)
+#if defined(__HIP__) || defined(__CUDA__)
+#define _DEFAULT_FN_ATTRS __attribute__((device))
+#else
+#define _DEFAULT_FN_ATTRS
+#endif
+#endif
+
 #if defined(__NVPTX__)
 #include <nvptxintrin.h>
 #elif defined(__AMDGPU__)
@@ -25,8 +33,15 @@
 #error "This header is only meant to be used on GPU architectures."
 #endif
 
+#if !defined(__cplusplus)
+#pragma push_macro("bool")
+#define bool _Bool
+#endif
+
+#pragma omp begin declare target device_type(nohost)
+
 // Returns the number of blocks in the requested dimension.
-_DEFAULT_FN_ATTRS static inline uint32_t __gpu_num_blocks(int __dim) {
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_blocks(int __dim) {
   switch (__dim) {
   case 0:
     return __gpu_num_blocks_x();
@@ -40,7 +55,7 @@ _DEFAULT_FN_ATTRS static inline uint32_t __gpu_num_blocks(int __dim) {
 }
 
 // Returns the number of block id in the requested dimension.
-_DEFAULT_FN_ATTRS static inline uint32_t __gpu_block_id(int __dim) {
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_block_id(int __dim) {
   switch (__dim) {
   case 0:
     return __gpu_block_id_x();
@@ -54,7 +69,7 @@ _DEFAULT_FN_ATTRS static inline uint32_t __gpu_block_id(int __dim) {
 }
 
 // Returns the number of threads in the requested dimension.
-_DEFAULT_FN_ATTRS static inline uint32_t __gpu_num_threads(int __dim) {
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_threads(int __dim) {
   switch (__dim) {
   case 0:
     return __gpu_num_threads_x();
@@ -68,7 +83,7 @@ _DEFAULT_FN_ATTRS static inline uint32_t __gpu_num_threads(int __dim) {
 }
 
 // Returns the thread id in the requested dimension.
-_DEFAULT_FN_ATTRS static inline uint32_t __gpu_thread_id(int __dim) {
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_thread_id(int __dim) {
   switch (__dim) {
   case 0:
     return __gpu_thread_id_x();
@@ -82,19 +97,19 @@ _DEFAULT_FN_ATTRS static inline uint32_t __gpu_thread_id(int __dim) {
 }
 
 // Get the first active thread inside the lane.
-_DEFAULT_FN_ATTRS static inline uint64_t
+_DEFAULT_FN_ATTRS static __inline__ uint64_t
 __gpu_first_lane_id(uint64_t __lane_mask) {
   return __builtin_ffsll(__lane_mask) - 1;
 }
 
 // Conditional that is only true for a single thread in a lane.
-_DEFAULT_FN_ATTRS static inline bool
+_DEFAULT_FN_ATTRS static __inline__ bool
 __gpu_is_first_in_lane(uint64_t __lane_mask) {
   return __gpu_lane_id() == __gpu_first_lane_id(__lane_mask);
 }
 
 // Gets the sum of all lanes inside the warp or wavefront.
-_DEFAULT_FN_ATTRS static inline uint32_t
+_DEFAULT_FN_ATTRS static __inline__ uint32_t
 __gpu_lane_reduce_u32(uint64_t __lane_mask, uint32_t x) {
   for (uint32_t step = __gpu_num_lanes() / 2; step > 0; step /= 2) {
     uint32_t index = step + __gpu_lane_id();
@@ -104,7 +119,7 @@ __gpu_lane_reduce_u32(uint64_t __lane_mask, uint32_t x) {
 }
 
 // Gets the accumulator scan of the threads in the warp or wavefront.
-_DEFAULT_FN_ATTRS static inline uint32_t
+_DEFAULT_FN_ATTRS static __inline__ uint32_t
 __gpu_lane_scan_u32(uint64_t __lane_mask, uint32_t x) {
   for (uint32_t step = 1; step < __gpu_num_lanes(); step *= 2) {
     uint32_t index = __gpu_lane_id() - step;
@@ -114,6 +129,12 @@ __gpu_lane_scan_u32(uint64_t __lane_mask, uint32_t x) {
   return x;
 }
 
+#pragma omp end declare target
+
+#if !defined(__cplusplus)
+#pragma pop_macro("bool")
+#endif
+
 #undef _DEFAULT_FN_ATTRS
 
 #endif // __GPUINTRIN_H