diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h b/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h index bcf01a2994..cd62c6626c 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h @@ -59,7 +59,7 @@ __device__ static inline int __clz(int input) { } __device__ static inline int __clzll(long long int input) { - return __ockl_clz_u64((__hip_uint64_t)input); + return __ockl_clz_u64((uint64_t)input); } __device__ static inline unsigned int __ffs(unsigned int input) { @@ -81,9 +81,9 @@ __device__ static inline unsigned int __ffsll(long long int input) { // Given a 32/64-bit value exec mask and an integer value base (between 0 and WAVEFRONT_SIZE), // find the n-th (given by offset) set bit in the exec mask from the base bit, and return the bit position. // If not found, return -1. -__device__ static __hip_int32_t __fns64(__hip_uint64_t mask, __hip_uint32_t base, __hip_int32_t offset) { - __hip_uint64_t temp_mask = mask; - __hip_int32_t temp_offset = offset; +__device__ static int32_t __fns64(uint64_t mask, uint32_t base, int32_t offset) { + uint64_t temp_mask = mask; + int32_t temp_offset = offset; if (offset == 0) { temp_mask &= (1 << base); @@ -98,10 +98,10 @@ __device__ static __hip_int32_t __fns64(__hip_uint64_t mask, __hip_uint32_t bas temp_mask = temp_mask & ((~0ULL) << base); if (__builtin_popcountll(temp_mask) < temp_offset) return -1; - __hip_int32_t total = 0; + int32_t total = 0; for (int i = 0x20; i > 0; i >>= 1) { - __hip_uint64_t temp_mask_lo = temp_mask & ((1ULL << i) - 1); - __hip_int32_t pcnt = __builtin_popcountll(temp_mask_lo); + uint64_t temp_mask_lo = temp_mask & ((1ULL << i) - 1); + int32_t pcnt = __builtin_popcountll(temp_mask_lo); if (pcnt < temp_offset) { temp_mask = temp_mask >> i; temp_offset -= pcnt; @@ -117,9 +117,9 @@ __device__ static __hip_int32_t __fns64(__hip_uint64_t mask, __hip_uint32_t bas return total; } -__device__ static __hip_int32_t __fns32(__hip_uint64_t mask, __hip_uint32_t base, __hip_int32_t offset) { - __hip_uint64_t temp_mask = mask; - __hip_int32_t temp_offset = offset; +__device__ static int32_t __fns32(uint64_t mask, uint32_t base, int32_t offset) { + uint64_t temp_mask = mask; + int32_t temp_offset = offset; if (offset == 0) { temp_mask &= (1 << base); temp_offset = 1; @@ -132,10 +132,10 @@ __device__ static __hip_int32_t __fns32(__hip_uint64_t mask, __hip_uint32_t base temp_mask = temp_mask & ((~0ULL) << base); if (__builtin_popcountll(temp_mask) < temp_offset) return -1; - __hip_int32_t total = 0; + int32_t total = 0; for (int i = 0x20; i > 0; i >>= 1) { - __hip_uint64_t temp_mask_lo = temp_mask & ((1ULL << i) - 1); - __hip_int32_t pcnt = __builtin_popcountll(temp_mask_lo); + uint64_t temp_mask_lo = temp_mask & ((1ULL << i) - 1); + int32_t pcnt = __builtin_popcountll(temp_mask_lo); if (pcnt < temp_offset) { temp_mask = temp_mask >> i; temp_offset -= pcnt; @@ -158,45 +158,45 @@ __device__ static inline unsigned long long int __brevll(unsigned long long int return __builtin_bitreverse64(input); } -__device__ static inline unsigned int __lastbit_u32_u64(__hip_uint64_t input) { +__device__ static inline unsigned int __lastbit_u32_u64(uint64_t input) { return input == 0 ? -1 : __builtin_ctzl(input); } __device__ static inline unsigned int __bitextract_u32(unsigned int src0, unsigned int src1, unsigned int src2) { - __hip_uint32_t offset = src1 & 31; - __hip_uint32_t width = src2 & 31; + uint32_t offset = src1 & 31; + uint32_t width = src2 & 31; return width == 0 ? 0 : (src0 << (32 - offset - width)) >> (32 - width); } -__device__ static inline __hip_uint64_t __bitextract_u64(__hip_uint64_t src0, unsigned int src1, unsigned int src2) { - __hip_uint64_t offset = src1 & 63; - __hip_uint64_t width = src2 & 63; +__device__ static inline uint64_t __bitextract_u64(uint64_t src0, unsigned int src1, unsigned int src2) { + uint64_t offset = src1 & 63; + uint64_t width = src2 & 63; return width == 0 ? 0 : (src0 << (64 - offset - width)) >> (64 - width); } __device__ static inline unsigned int __bitinsert_u32(unsigned int src0, unsigned int src1, unsigned int src2, unsigned int src3) { - __hip_uint32_t offset = src2 & 31; - __hip_uint32_t width = src3 & 31; - __hip_uint32_t mask = (1 << width) - 1; + uint32_t offset = src2 & 31; + uint32_t width = src3 & 31; + uint32_t mask = (1 << width) - 1; return ((src0 & ~(mask << offset)) | ((src1 & mask) << offset)); } -__device__ static inline __hip_uint64_t __bitinsert_u64(__hip_uint64_t src0, __hip_uint64_t src1, unsigned int src2, unsigned int src3) { - __hip_uint64_t offset = src2 & 63; - __hip_uint64_t width = src3 & 63; - __hip_uint64_t mask = (1ULL << width) - 1; +__device__ static inline uint64_t __bitinsert_u64(uint64_t src0, uint64_t src1, unsigned int src2, unsigned int src3) { + uint64_t offset = src2 & 63; + uint64_t width = src3 & 63; + uint64_t mask = (1ULL << width) - 1; return ((src0 & ~(mask << offset)) | ((src1 & mask) << offset)); } __device__ inline unsigned int __funnelshift_l(unsigned int lo, unsigned int hi, unsigned int shift) { - __hip_uint32_t mask_shift = shift & 31; + uint32_t mask_shift = shift & 31; return mask_shift == 0 ? hi : __builtin_amdgcn_alignbit(hi, lo, 32 - mask_shift); } __device__ inline unsigned int __funnelshift_lc(unsigned int lo, unsigned int hi, unsigned int shift) { - __hip_uint32_t min_shift = shift >= 32 ? 32 : shift; + uint32_t min_shift = shift >= 32 ? 32 : shift; return min_shift == 0 ? hi : __builtin_amdgcn_alignbit(hi, lo, 32 - min_shift); } @@ -525,9 +525,9 @@ __device__ static inline unsigned int __float_as_uint(float x) { } __device__ static inline double __hiloint2double(int hi, int lo) { - static_assert(sizeof(double) == sizeof(__hip_uint64_t), ""); + static_assert(sizeof(double) == sizeof(uint64_t), ""); - __hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(hi) << 32ull) | static_cast<__hip_uint32_t>(lo); + uint64_t tmp0 = (static_cast(hi) << 32ull) | static_cast(lo); double tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); @@ -717,32 +717,32 @@ unsigned long long int __ballot64(int predicate) { // hip.amdgcn.bc - lanemask __device__ inline -__hip_uint64_t __lanemask_gt() +uint64_t __lanemask_gt() { - __hip_uint32_t lane = __ockl_lane_u32(); + uint32_t lane = __ockl_lane_u32(); if (lane == 63) return 0; - __hip_uint64_t ballot = __ballot64(1); - __hip_uint64_t mask = (~((__hip_uint64_t)0)) << (lane + 1); + uint64_t ballot = __ballot64(1); + uint64_t mask = (~((uint64_t)0)) << (lane + 1); return mask & ballot; } __device__ inline -__hip_uint64_t __lanemask_lt() +uint64_t __lanemask_lt() { - __hip_uint32_t lane = __ockl_lane_u32(); - __hip_int64_t ballot = __ballot64(1); - __hip_uint64_t mask = ((__hip_uint64_t)1 << lane) - (__hip_uint64_t)1; + uint32_t lane = __ockl_lane_u32(); + int64_t ballot = __ballot64(1); + uint64_t mask = ((uint64_t)1 << lane) - (uint64_t)1; return mask & ballot; } __device__ inline -__hip_uint64_t __lanemask_eq() +uint64_t __lanemask_eq() { - __hip_uint32_t lane = __ockl_lane_u32(); - __hip_int64_t mask = ((__hip_uint64_t)1 << lane); + uint32_t lane = __ockl_lane_u32(); + int64_t mask = ((uint64_t)1 << lane); return mask; } diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_bfloat16.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_bfloat16.h index 937486f149..deb3bfb7e2 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_bfloat16.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_bfloat16.h @@ -81,9 +81,9 @@ struct hip_bfloat16 { union { - __hip_uint32_t int32; + uint32_t int32; float fp32; - } u = {__hip_uint32_t(data) << 16}; + } u = {uint32_t(data) << 16}; return u.fp32; } @@ -113,7 +113,7 @@ private: union { float fp32; - __hip_uint32_t int32; + uint32_t int32; } u = {f}; if(~u.int32 & 0x7f800000) { @@ -156,7 +156,7 @@ private: union { float fp32; - __hip_uint32_t int32; + uint32_t int32; } u = {f}; return __hip_uint16_t(u.int32 >> 16) | (!(~u.int32 & 0x7f800000) && (u.int32 & 0xffff)); } diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h index 7b8e2e5b57..eeb67bd075 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h @@ -62,9 +62,9 @@ namespace cooperative_groups { */ class thread_group { protected: - __hip_uint32_t _type; // thread_group type - __hip_uint32_t _size; // total number of threads in the tread_group - __hip_uint64_t _mask; // Lanemask for coalesced and tiled partitioned group types, + uint32_t _type; // thread_group type + uint32_t _size; // total number of threads in the tread_group + uint64_t _mask; // Lanemask for coalesced and tiled partitioned group types, // LSB represents lane 0, and MSB represents lane 63 // Construct a thread group, and set thread group type and other essential @@ -72,8 +72,8 @@ class thread_group { // only when the group is supposed to contain only the calling the thread // (throurh the API - `this_thread()`), and in all other cases, this thread // group object is a sub-object of some other derived thread group object - __CG_QUALIFIER__ thread_group(internal::group_type type, __hip_uint32_t size = static_cast<__hip_uint64_t>(0), - __hip_uint64_t mask = static_cast<__hip_uint64_t>(0)) { + __CG_QUALIFIER__ thread_group(internal::group_type type, uint32_t size = static_cast(0), + uint64_t mask = static_cast(0)) { _type = type; _size = size; _mask = mask; @@ -100,10 +100,10 @@ class thread_group { // Total number of threads in the thread group, and this serves the purpose // for all derived cooperative group types since their `size` is directly // saved during the construction - __CG_QUALIFIER__ __hip_uint32_t size() const { return _size; } + __CG_QUALIFIER__ uint32_t size() const { return _size; } __CG_QUALIFIER__ unsigned int cg_type() const { return _type; } // Rank of the calling thread within [0, size()) - __CG_QUALIFIER__ __hip_uint32_t thread_rank() const; + __CG_QUALIFIER__ uint32_t thread_rank() const; // Is this cooperative group type valid? __CG_QUALIFIER__ bool is_valid() const; // synchronize the threads in the thread group @@ -139,17 +139,17 @@ class multi_grid_group : public thread_group { protected: // Construct mutli-grid thread group (through the API this_multi_grid()) - explicit __CG_QUALIFIER__ multi_grid_group(__hip_uint32_t size) + explicit __CG_QUALIFIER__ multi_grid_group(uint32_t size) : thread_group(internal::cg_multi_grid, size) {} public: // Number of invocations participating in this multi-grid group. In other // words, the number of GPUs - __CG_QUALIFIER__ __hip_uint32_t num_grids() { return internal::multi_grid::num_grids(); } + __CG_QUALIFIER__ uint32_t num_grids() { return internal::multi_grid::num_grids(); } // Rank of this invocation. In other words, an ID number within the range // [0, num_grids()) of the GPU, this kernel is running on - __CG_QUALIFIER__ __hip_uint32_t grid_rank() { return internal::multi_grid::grid_rank(); } - __CG_QUALIFIER__ __hip_uint32_t thread_rank() const { return internal::multi_grid::thread_rank(); } + __CG_QUALIFIER__ uint32_t grid_rank() { return internal::multi_grid::grid_rank(); } + __CG_QUALIFIER__ uint32_t thread_rank() const { return internal::multi_grid::thread_rank(); } __CG_QUALIFIER__ bool is_valid() const { return internal::multi_grid::is_valid(); } __CG_QUALIFIER__ void sync() const { internal::multi_grid::sync(); } }; @@ -182,10 +182,10 @@ class grid_group : public thread_group { protected: // Construct grid thread group (through the API this_grid()) - explicit __CG_QUALIFIER__ grid_group(__hip_uint32_t size) : thread_group(internal::cg_grid, size) {} + explicit __CG_QUALIFIER__ grid_group(uint32_t size) : thread_group(internal::cg_grid, size) {} public: - __CG_QUALIFIER__ __hip_uint32_t thread_rank() const { return internal::grid::thread_rank(); } + __CG_QUALIFIER__ uint32_t thread_rank() const { return internal::grid::thread_rank(); } __CG_QUALIFIER__ bool is_valid() const { return internal::grid::is_valid(); } __CG_QUALIFIER__ void sync() const { internal::grid::sync(); } }; @@ -220,7 +220,7 @@ class thread_block : public thread_group { unsigned int tile_size); protected: // Construct a workgroup thread group (through the API this_thread_block()) - explicit __CG_QUALIFIER__ thread_block(__hip_uint32_t size) + explicit __CG_QUALIFIER__ thread_block(uint32_t size) : thread_group(internal::cg_workgroup, size) {} __CG_QUALIFIER__ thread_group new_tiled_group(unsigned int tile_size) const { @@ -243,8 +243,8 @@ class thread_block : public thread_group { __CG_STATIC_QUALIFIER__ dim3 group_index() { return internal::workgroup::group_index(); } // 3-dimensional thread index within the block __CG_STATIC_QUALIFIER__ dim3 thread_index() { return internal::workgroup::thread_index(); } - __CG_STATIC_QUALIFIER__ __hip_uint32_t thread_rank() { return internal::workgroup::thread_rank(); } - __CG_STATIC_QUALIFIER__ __hip_uint32_t size() { return internal::workgroup::size(); } + __CG_STATIC_QUALIFIER__ uint32_t thread_rank() { return internal::workgroup::thread_rank(); } + __CG_STATIC_QUALIFIER__ uint32_t size() { return internal::workgroup::size(); } __CG_STATIC_QUALIFIER__ bool is_valid() { return internal::workgroup::is_valid(); } __CG_STATIC_QUALIFIER__ void sync() { internal::workgroup::sync(); } __CG_QUALIFIER__ dim3 group_dim() { return internal::workgroup::block_dim(); } @@ -490,7 +490,7 @@ __CG_QUALIFIER__ coalesced_group coalesced_threads() { * @note This function is implemented on Linux, under developement * on Windows. */ -__CG_QUALIFIER__ __hip_uint32_t thread_group::thread_rank() const { +__CG_QUALIFIER__ uint32_t thread_group::thread_rank() const { switch (this->_type) { case internal::cg_multi_grid: { return (static_cast(this)->thread_rank()); @@ -580,14 +580,14 @@ __CG_QUALIFIER__ void thread_group::sync() const { * @note This function is implemented on Linux, under developement * on Windows. */ -template __CG_QUALIFIER__ __hip_uint32_t group_size(CGTy const& g) { return g.size(); } +template __CG_QUALIFIER__ uint32_t group_size(CGTy const& g) { return g.size(); } /** * Implemenation of publicly exposed `wrapper` API on top of basic cooperative * group type APIs * @note This function is implemented on Linux, under developement * on Windows. */ -template __CG_QUALIFIER__ __hip_uint32_t thread_rank(CGTy const& g) { +template __CG_QUALIFIER__ uint32_t thread_rank(CGTy const& g) { return g.thread_rank(); } /** diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime.h index 0d8f507534..903de73fb1 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime.h @@ -87,6 +87,15 @@ size_t amd_dbgapi_get_build_id(); #include #include #endif // __cplusplus +#else +#if !__HIP_NO_STD_DEFS__ +typedef unsigned int uint32_t; +typedef signed int int32_t; +namespace std { +using ::uint32_t; +using ::int32_t; +} +#endif // __HIP_NO_STD_DEFS__ #endif // !defined(__HIPCC_RTC__) #if __HIP_CLANG_ONLY__ @@ -240,35 +249,34 @@ extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_size(unsign extern "C" __device__ __attribute__((const)) size_t __ockl_get_num_groups(unsigned int); struct __HIP_BlockIdx { __device__ - __hip_uint32_t operator()(__hip_uint32_t x) const noexcept { return __ockl_get_group_id(x); } + std::uint32_t operator()(std::uint32_t x) const noexcept { return __ockl_get_group_id(x); } }; struct __HIP_BlockDim { __device__ - __hip_uint32_t operator()(__hip_uint32_t x) const noexcept { + std::uint32_t operator()(std::uint32_t x) const noexcept { return __ockl_get_local_size(x); } }; struct __HIP_GridDim { __device__ - __hip_uint32_t operator()(__hip_uint32_t x) const noexcept { + std::uint32_t operator()(std::uint32_t x) const noexcept { return __ockl_get_num_groups(x); } }; struct __HIP_ThreadIdx { __device__ - __hip_uint32_t operator()(__hip_uint32_t x) const noexcept { + std::uint32_t operator()(std::uint32_t x) const noexcept { return __ockl_get_local_id(x); } }; #if defined(__HIPCC_RTC__) typedef struct dim3 { - __hip_uint32_t x; ///< x - __hip_uint32_t y; ///< y - __hip_uint32_t z; ///< z + uint32_t x; ///< x + uint32_t y; ///< y + uint32_t z; ///< z #ifdef __cplusplus - constexpr __device__ dim3(__hip_uint32_t _x = 1, __hip_uint32_t _y = 1, __hip_uint32_t _z = 1) - : x(_x), y(_y), z(_z){}; + constexpr __device__ dim3(uint32_t _x = 1, uint32_t _y = 1, uint32_t _z = 1) : x(_x), y(_y), z(_z){}; #endif } dim3; #endif // !defined(__HIPCC_RTC__) @@ -317,37 +325,37 @@ constexpr typename __HIP_Coordinates::__Z __HIP_Coordinates::z; inline __device__ -__hip_uint32_t operator*(__HIP_Coordinates<__HIP_GridDim>::__X, +std::uint32_t operator*(__HIP_Coordinates<__HIP_GridDim>::__X, __HIP_Coordinates<__HIP_BlockDim>::__X) noexcept { return __ockl_get_global_size(0); } inline __device__ -__hip_uint32_t operator*(__HIP_Coordinates<__HIP_BlockDim>::__X, +std::uint32_t operator*(__HIP_Coordinates<__HIP_BlockDim>::__X, __HIP_Coordinates<__HIP_GridDim>::__X) noexcept { return __ockl_get_global_size(0); } inline __device__ -__hip_uint32_t operator*(__HIP_Coordinates<__HIP_GridDim>::__Y, +std::uint32_t operator*(__HIP_Coordinates<__HIP_GridDim>::__Y, __HIP_Coordinates<__HIP_BlockDim>::__Y) noexcept { return __ockl_get_global_size(1); } inline __device__ -__hip_uint32_t operator*(__HIP_Coordinates<__HIP_BlockDim>::__Y, +std::uint32_t operator*(__HIP_Coordinates<__HIP_BlockDim>::__Y, __HIP_Coordinates<__HIP_GridDim>::__Y) noexcept { return __ockl_get_global_size(1); } inline __device__ -__hip_uint32_t operator*(__HIP_Coordinates<__HIP_GridDim>::__Z, +std::uint32_t operator*(__HIP_Coordinates<__HIP_GridDim>::__Z, __HIP_Coordinates<__HIP_BlockDim>::__Z) noexcept { return __ockl_get_global_size(2); } inline __device__ -__hip_uint32_t operator*(__HIP_Coordinates<__HIP_BlockDim>::__Z, +std::uint32_t operator*(__HIP_Coordinates<__HIP_BlockDim>::__Z, __HIP_Coordinates<__HIP_GridDim>::__Z) noexcept { return __ockl_get_global_size(2); } diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_warp_functions.h b/projects/clr/hipamd/include/hip/amd_detail/amd_warp_functions.h index 36bd23af2f..559ab20b33 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_warp_functions.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_warp_functions.h @@ -100,13 +100,13 @@ __device__ inline double __shfl(double var, int src_lane, int width = warpSize) { static_assert(sizeof(double) == 2 * sizeof(int), ""); - static_assert(sizeof(double) == sizeof(__hip_uint64_t), ""); + static_assert(sizeof(double) == sizeof(uint64_t), ""); int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); tmp[0] = __shfl(tmp[0], src_lane, width); tmp[1] = __shfl(tmp[1], src_lane, width); - __hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(tmp[0]); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); double tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); return tmp1; } @@ -116,13 +116,13 @@ long __shfl(long var, int src_lane, int width = warpSize) { #ifndef _MSC_VER static_assert(sizeof(long) == 2 * sizeof(int), ""); - static_assert(sizeof(long) == sizeof(__hip_uint64_t), ""); + static_assert(sizeof(long) == sizeof(uint64_t), ""); int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); tmp[0] = __shfl(tmp[0], src_lane, width); tmp[1] = __shfl(tmp[1], src_lane, width); - __hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(tmp[0]); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); return tmp1; #else @@ -135,13 +135,13 @@ inline unsigned long __shfl(unsigned long var, int src_lane, int width = warpSize) { #ifndef _MSC_VER static_assert(sizeof(unsigned long) == 2 * sizeof(unsigned int), ""); - static_assert(sizeof(unsigned long) == sizeof(__hip_uint64_t), ""); + static_assert(sizeof(unsigned long) == sizeof(uint64_t), ""); unsigned int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); tmp[0] = __shfl(tmp[0], src_lane, width); tmp[1] = __shfl(tmp[1], src_lane, width); - __hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(tmp[0]); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); unsigned long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); return tmp1; #else @@ -154,13 +154,13 @@ inline long long __shfl(long long var, int src_lane, int width = warpSize) { static_assert(sizeof(long long) == 2 * sizeof(int), ""); - static_assert(sizeof(long long) == sizeof(__hip_uint64_t), ""); + static_assert(sizeof(long long) == sizeof(uint64_t), ""); int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); tmp[0] = __shfl(tmp[0], src_lane, width); tmp[1] = __shfl(tmp[1], src_lane, width); - __hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(tmp[0]); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); return tmp1; } @@ -168,13 +168,13 @@ __device__ inline unsigned long long __shfl(unsigned long long var, int src_lane, int width = warpSize) { static_assert(sizeof(unsigned long long) == 2 * sizeof(unsigned int), ""); - static_assert(sizeof(unsigned long long) == sizeof(__hip_uint64_t), ""); + static_assert(sizeof(unsigned long long) == sizeof(uint64_t), ""); unsigned int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); tmp[0] = __shfl(tmp[0], src_lane, width); tmp[1] = __shfl(tmp[1], src_lane, width); - __hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(tmp[0]); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); unsigned long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); return tmp1; } @@ -205,13 +205,13 @@ __device__ inline double __shfl_up(double var, unsigned int lane_delta, int width = warpSize) { static_assert(sizeof(double) == 2 * sizeof(int), ""); - static_assert(sizeof(double) == sizeof(__hip_uint64_t), ""); + static_assert(sizeof(double) == sizeof(uint64_t), ""); int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); tmp[0] = __shfl_up(tmp[0], lane_delta, width); tmp[1] = __shfl_up(tmp[1], lane_delta, width); - __hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(tmp[0]); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); double tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); return tmp1; } @@ -221,13 +221,13 @@ long __shfl_up(long var, unsigned int lane_delta, int width = warpSize) { #ifndef _MSC_VER static_assert(sizeof(long) == 2 * sizeof(int), ""); - static_assert(sizeof(long) == sizeof(__hip_uint64_t), ""); + static_assert(sizeof(long) == sizeof(uint64_t), ""); int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); tmp[0] = __shfl_up(tmp[0], lane_delta, width); tmp[1] = __shfl_up(tmp[1], lane_delta, width); - __hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(tmp[0]); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); return tmp1; #else @@ -242,13 +242,13 @@ unsigned long __shfl_up(unsigned long var, unsigned int lane_delta, int width = { #ifndef _MSC_VER static_assert(sizeof(unsigned long) == 2 * sizeof(unsigned int), ""); - static_assert(sizeof(unsigned long) == sizeof(__hip_uint64_t), ""); + static_assert(sizeof(unsigned long) == sizeof(uint64_t), ""); unsigned int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); tmp[0] = __shfl_up(tmp[0], lane_delta, width); tmp[1] = __shfl_up(tmp[1], lane_delta, width); - __hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(tmp[0]); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); unsigned long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); return tmp1; #else @@ -262,11 +262,11 @@ inline long long __shfl_up(long long var, unsigned int lane_delta, int width = warpSize) { static_assert(sizeof(long long) == 2 * sizeof(int), ""); - static_assert(sizeof(long long) == sizeof(__hip_uint64_t), ""); + static_assert(sizeof(long long) == sizeof(uint64_t), ""); int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); tmp[0] = __shfl_up(tmp[0], lane_delta, width); tmp[1] = __shfl_up(tmp[1], lane_delta, width); - __hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(tmp[0]); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); return tmp1; } @@ -276,11 +276,11 @@ inline unsigned long long __shfl_up(unsigned long long var, unsigned int lane_delta, int width = warpSize) { static_assert(sizeof(unsigned long long) == 2 * sizeof(unsigned int), ""); - static_assert(sizeof(unsigned long long) == sizeof(__hip_uint64_t), ""); + static_assert(sizeof(unsigned long long) == sizeof(uint64_t), ""); unsigned int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); tmp[0] = __shfl_up(tmp[0], lane_delta, width); tmp[1] = __shfl_up(tmp[1], lane_delta, width); - __hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(tmp[0]); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); unsigned long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); return tmp1; } @@ -311,13 +311,13 @@ __device__ inline double __shfl_down(double var, unsigned int lane_delta, int width = warpSize) { static_assert(sizeof(double) == 2 * sizeof(int), ""); - static_assert(sizeof(double) == sizeof(__hip_uint64_t), ""); + static_assert(sizeof(double) == sizeof(uint64_t), ""); int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); tmp[0] = __shfl_down(tmp[0], lane_delta, width); tmp[1] = __shfl_down(tmp[1], lane_delta, width); - __hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(tmp[0]); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); double tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); return tmp1; } @@ -327,13 +327,13 @@ long __shfl_down(long var, unsigned int lane_delta, int width = warpSize) { #ifndef _MSC_VER static_assert(sizeof(long) == 2 * sizeof(int), ""); - static_assert(sizeof(long) == sizeof(__hip_uint64_t), ""); + static_assert(sizeof(long) == sizeof(uint64_t), ""); int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); tmp[0] = __shfl_down(tmp[0], lane_delta, width); tmp[1] = __shfl_down(tmp[1], lane_delta, width); - __hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(tmp[0]); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); return tmp1; #else @@ -347,13 +347,13 @@ unsigned long __shfl_down(unsigned long var, unsigned int lane_delta, int width { #ifndef _MSC_VER static_assert(sizeof(unsigned long) == 2 * sizeof(unsigned int), ""); - static_assert(sizeof(unsigned long) == sizeof(__hip_uint64_t), ""); + static_assert(sizeof(unsigned long) == sizeof(uint64_t), ""); unsigned int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); tmp[0] = __shfl_down(tmp[0], lane_delta, width); tmp[1] = __shfl_down(tmp[1], lane_delta, width); - __hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(tmp[0]); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); unsigned long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); return tmp1; #else @@ -366,11 +366,11 @@ inline long long __shfl_down(long long var, unsigned int lane_delta, int width = warpSize) { static_assert(sizeof(long long) == 2 * sizeof(int), ""); - static_assert(sizeof(long long) == sizeof(__hip_uint64_t), ""); + static_assert(sizeof(long long) == sizeof(uint64_t), ""); int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); tmp[0] = __shfl_down(tmp[0], lane_delta, width); tmp[1] = __shfl_down(tmp[1], lane_delta, width); - __hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(tmp[0]); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); return tmp1; } @@ -379,11 +379,11 @@ inline unsigned long long __shfl_down(unsigned long long var, unsigned int lane_delta, int width = warpSize) { static_assert(sizeof(unsigned long long) == 2 * sizeof(unsigned int), ""); - static_assert(sizeof(unsigned long long) == sizeof(__hip_uint64_t), ""); + static_assert(sizeof(unsigned long long) == sizeof(uint64_t), ""); unsigned int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); tmp[0] = __shfl_down(tmp[0], lane_delta, width); tmp[1] = __shfl_down(tmp[1], lane_delta, width); - __hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(tmp[0]); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); unsigned long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); return tmp1; } @@ -414,13 +414,13 @@ __device__ inline double __shfl_xor(double var, int lane_mask, int width = warpSize) { static_assert(sizeof(double) == 2 * sizeof(int), ""); - static_assert(sizeof(double) == sizeof(__hip_uint64_t), ""); + static_assert(sizeof(double) == sizeof(uint64_t), ""); int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); tmp[0] = __shfl_xor(tmp[0], lane_mask, width); tmp[1] = __shfl_xor(tmp[1], lane_mask, width); - __hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(tmp[0]); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); double tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); return tmp1; } @@ -430,13 +430,13 @@ long __shfl_xor(long var, int lane_mask, int width = warpSize) { #ifndef _MSC_VER static_assert(sizeof(long) == 2 * sizeof(int), ""); - static_assert(sizeof(long) == sizeof(__hip_uint64_t), ""); + static_assert(sizeof(long) == sizeof(uint64_t), ""); int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); tmp[0] = __shfl_xor(tmp[0], lane_mask, width); tmp[1] = __shfl_xor(tmp[1], lane_mask, width); - __hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(tmp[0]); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); return tmp1; #else @@ -450,13 +450,13 @@ unsigned long __shfl_xor(unsigned long var, int lane_mask, int width = warpSize) { #ifndef _MSC_VER static_assert(sizeof(unsigned long) == 2 * sizeof(unsigned int), ""); - static_assert(sizeof(unsigned long) == sizeof(__hip_uint64_t), ""); + static_assert(sizeof(unsigned long) == sizeof(uint64_t), ""); unsigned int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); tmp[0] = __shfl_xor(tmp[0], lane_mask, width); tmp[1] = __shfl_xor(tmp[1], lane_mask, width); - __hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(tmp[0]); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); unsigned long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); return tmp1; #else @@ -469,11 +469,11 @@ inline long long __shfl_xor(long long var, int lane_mask, int width = warpSize) { static_assert(sizeof(long long) == 2 * sizeof(int), ""); - static_assert(sizeof(long long) == sizeof(__hip_uint64_t), ""); + static_assert(sizeof(long long) == sizeof(uint64_t), ""); int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); tmp[0] = __shfl_xor(tmp[0], lane_mask, width); tmp[1] = __shfl_xor(tmp[1], lane_mask, width); - __hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(tmp[0]); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); return tmp1; } @@ -482,11 +482,11 @@ inline unsigned long long __shfl_xor(unsigned long long var, int lane_mask, int width = warpSize) { static_assert(sizeof(unsigned long long) == 2 * sizeof(unsigned int), ""); - static_assert(sizeof(unsigned long long) == sizeof(__hip_uint64_t), ""); + static_assert(sizeof(unsigned long long) == sizeof(uint64_t), ""); unsigned int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); tmp[0] = __shfl_xor(tmp[0], lane_mask, width); tmp[1] = __shfl_xor(tmp[1], lane_mask, width); - __hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(tmp[0]); + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); unsigned long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); return tmp1; } diff --git a/projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h b/projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h index 9b543df47f..edc4692c83 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h +++ b/projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h @@ -54,7 +54,7 @@ extern "C" __device__ __attribute__((const)) uint __ockl_sadd_u32(uint, uint, ui extern "C" __device__ __attribute__((const)) uchar __ockl_clz_u8(uchar); extern "C" __device__ __attribute__((const)) ushort __ockl_clz_u16(ushort); extern "C" __device__ __attribute__((const)) uint __ockl_clz_u32(uint); -extern "C" __device__ __attribute__((const)) __hip_uint64_t __ockl_clz_u64(__hip_uint64_t); +extern "C" __device__ __attribute__((const)) uint64_t __ockl_clz_u64(uint64_t); extern "C" __device__ __attribute__((const)) float __ocml_floor_f32(float); extern "C" __device__ __attribute__((const)) float __ocml_rint_f32(float); @@ -75,26 +75,26 @@ extern "C" __device__ __attribute__((const)) _Float16 __ocml_cvtrtz_f16_f32(floa extern "C" __device__ __attribute__((const)) float __ocml_cvtrtn_f32_s32(int); extern "C" __device__ __attribute__((const)) float __ocml_cvtrtp_f32_s32(int); extern "C" __device__ __attribute__((const)) float __ocml_cvtrtz_f32_s32(int); -extern "C" __device__ __attribute__((const)) float __ocml_cvtrtn_f32_u32(__hip_uint32_t); -extern "C" __device__ __attribute__((const)) float __ocml_cvtrtp_f32_u32(__hip_uint32_t); -extern "C" __device__ __attribute__((const)) float __ocml_cvtrtz_f32_u32(__hip_uint32_t); -extern "C" __device__ __attribute__((const)) float __ocml_cvtrtn_f32_s64(__hip_int64_t); -extern "C" __device__ __attribute__((const)) float __ocml_cvtrtp_f32_s64(__hip_int64_t); -extern "C" __device__ __attribute__((const)) float __ocml_cvtrtz_f32_s64(__hip_int64_t); -extern "C" __device__ __attribute__((const)) float __ocml_cvtrtn_f32_u64(__hip_uint64_t); -extern "C" __device__ __attribute__((const)) float __ocml_cvtrtp_f32_u64(__hip_uint64_t); -extern "C" __device__ __attribute__((const)) float __ocml_cvtrtz_f32_u64(__hip_uint64_t); -extern "C" __device__ __attribute__((const)) double __ocml_cvtrtn_f64_s64(__hip_int64_t); -extern "C" __device__ __attribute__((const)) double __ocml_cvtrtp_f64_s64(__hip_int64_t); -extern "C" __device__ __attribute__((const)) double __ocml_cvtrtz_f64_s64(__hip_int64_t); -extern "C" __device__ __attribute__((const)) double __ocml_cvtrtn_f64_u64(__hip_uint64_t); -extern "C" __device__ __attribute__((const)) double __ocml_cvtrtp_f64_u64(__hip_uint64_t); -extern "C" __device__ __attribute__((const)) double __ocml_cvtrtz_f64_u64(__hip_uint64_t); +extern "C" __device__ __attribute__((const)) float __ocml_cvtrtn_f32_u32(uint32_t); +extern "C" __device__ __attribute__((const)) float __ocml_cvtrtp_f32_u32(uint32_t); +extern "C" __device__ __attribute__((const)) float __ocml_cvtrtz_f32_u32(uint32_t); +extern "C" __device__ __attribute__((const)) float __ocml_cvtrtn_f32_s64(int64_t); +extern "C" __device__ __attribute__((const)) float __ocml_cvtrtp_f32_s64(int64_t); +extern "C" __device__ __attribute__((const)) float __ocml_cvtrtz_f32_s64(int64_t); +extern "C" __device__ __attribute__((const)) float __ocml_cvtrtn_f32_u64(uint64_t); +extern "C" __device__ __attribute__((const)) float __ocml_cvtrtp_f32_u64(uint64_t); +extern "C" __device__ __attribute__((const)) float __ocml_cvtrtz_f32_u64(uint64_t); +extern "C" __device__ __attribute__((const)) double __ocml_cvtrtn_f64_s64(int64_t); +extern "C" __device__ __attribute__((const)) double __ocml_cvtrtp_f64_s64(int64_t); +extern "C" __device__ __attribute__((const)) double __ocml_cvtrtz_f64_s64(int64_t); +extern "C" __device__ __attribute__((const)) double __ocml_cvtrtn_f64_u64(uint64_t); +extern "C" __device__ __attribute__((const)) double __ocml_cvtrtp_f64_u64(uint64_t); +extern "C" __device__ __attribute__((const)) double __ocml_cvtrtz_f64_u64(uint64_t); extern "C" __device__ __attribute__((convergent)) void __ockl_gws_init(uint nwm1, uint rid); extern "C" __device__ __attribute__((convergent)) void __ockl_gws_barrier(uint nwm1, uint rid); -extern "C" __device__ __attribute__((const)) __hip_uint32_t __ockl_lane_u32(); +extern "C" __device__ __attribute__((const)) uint32_t __ockl_lane_u32(); extern "C" __device__ __attribute__((const)) int __ockl_grid_is_valid(void); extern "C" __device__ __attribute__((convergent)) void __ockl_grid_sync(void); extern "C" __device__ __attribute__((const)) uint __ockl_multi_grid_num_grids(void); @@ -110,14 +110,14 @@ extern "C" __device__ __attribute__((convergent)) int __ockl_wgred_add_i32(int a extern "C" __device__ __attribute__((convergent)) int __ockl_wgred_and_i32(int a); extern "C" __device__ __attribute__((convergent)) int __ockl_wgred_or_i32(int a); -extern "C" __device__ __hip_uint64_t __ockl_fprintf_stderr_begin(); -extern "C" __device__ __hip_uint64_t __ockl_fprintf_append_args(__hip_uint64_t msg_desc, __hip_uint32_t num_args, - __hip_uint64_t value0, __hip_uint64_t value1, - __hip_uint64_t value2, __hip_uint64_t value3, - __hip_uint64_t value4, __hip_uint64_t value5, - __hip_uint64_t value6, __hip_uint32_t is_last); -extern "C" __device__ __hip_uint64_t __ockl_fprintf_append_string_n(__hip_uint64_t msg_desc, const char* data, - __hip_uint64_t length, __hip_uint32_t is_last); +extern "C" __device__ uint64_t __ockl_fprintf_stderr_begin(); +extern "C" __device__ uint64_t __ockl_fprintf_append_args(uint64_t msg_desc, uint32_t num_args, + uint64_t value0, uint64_t value1, + uint64_t value2, uint64_t value3, + uint64_t value4, uint64_t value5, + uint64_t value6, uint32_t is_last); +extern "C" __device__ uint64_t __ockl_fprintf_append_string_n(uint64_t msg_desc, const char* data, + uint64_t length, uint32_t is_last); // Introduce local address space #define __local __attribute__((address_space(3))) diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h index 880e6207da..15607b9683 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h @@ -112,16 +112,16 @@ typedef enum { */ namespace multi_grid { -__CG_STATIC_QUALIFIER__ __hip_uint32_t num_grids() { - return static_cast<__hip_uint32_t>(__ockl_multi_grid_num_grids()); } +__CG_STATIC_QUALIFIER__ uint32_t num_grids() { + return static_cast(__ockl_multi_grid_num_grids()); } -__CG_STATIC_QUALIFIER__ __hip_uint32_t grid_rank() { - return static_cast<__hip_uint32_t>(__ockl_multi_grid_grid_rank()); } +__CG_STATIC_QUALIFIER__ uint32_t grid_rank() { + return static_cast(__ockl_multi_grid_grid_rank()); } -__CG_STATIC_QUALIFIER__ __hip_uint32_t size() { return static_cast<__hip_uint32_t>(__ockl_multi_grid_size()); } +__CG_STATIC_QUALIFIER__ uint32_t size() { return static_cast(__ockl_multi_grid_size()); } -__CG_STATIC_QUALIFIER__ __hip_uint32_t thread_rank() { - return static_cast<__hip_uint32_t>(__ockl_multi_grid_thread_rank()); } +__CG_STATIC_QUALIFIER__ uint32_t thread_rank() { + return static_cast(__ockl_multi_grid_thread_rank()); } __CG_STATIC_QUALIFIER__ bool is_valid() { return static_cast(__ockl_multi_grid_is_valid()); } @@ -135,23 +135,23 @@ __CG_STATIC_QUALIFIER__ void sync() { __ockl_multi_grid_sync(); } */ namespace grid { -__CG_STATIC_QUALIFIER__ __hip_uint32_t size() { - return static_cast<__hip_uint32_t>((blockDim.z * gridDim.z) * (blockDim.y * gridDim.y) * +__CG_STATIC_QUALIFIER__ uint32_t size() { + return static_cast((blockDim.z * gridDim.z) * (blockDim.y * gridDim.y) * (blockDim.x * gridDim.x)); } -__CG_STATIC_QUALIFIER__ __hip_uint32_t thread_rank() { +__CG_STATIC_QUALIFIER__ uint32_t thread_rank() { // Compute global id of the workgroup to which the current thread belongs to - __hip_uint32_t blkIdx = static_cast<__hip_uint32_t>((blockIdx.z * gridDim.y * gridDim.x) + + uint32_t blkIdx = static_cast((blockIdx.z * gridDim.y * gridDim.x) + (blockIdx.y * gridDim.x) + (blockIdx.x)); // Compute total number of threads being passed to reach current workgroup // within grid - __hip_uint32_t num_threads_till_current_workgroup = - static_cast<__hip_uint32_t>(blkIdx * (blockDim.x * blockDim.y * blockDim.z)); + uint32_t num_threads_till_current_workgroup = + static_cast(blkIdx * (blockDim.x * blockDim.y * blockDim.z)); // Compute thread local rank within current workgroup - __hip_uint32_t local_thread_rank = static_cast<__hip_uint32_t>((threadIdx.z * blockDim.y * blockDim.x) + + uint32_t local_thread_rank = static_cast((threadIdx.z * blockDim.y * blockDim.x) + (threadIdx.y * blockDim.x) + (threadIdx.x)); return (num_threads_till_current_workgroup + local_thread_rank); @@ -171,21 +171,21 @@ __CG_STATIC_QUALIFIER__ void sync() { __ockl_grid_sync(); } namespace workgroup { __CG_STATIC_QUALIFIER__ dim3 group_index() { - return (dim3(static_cast<__hip_uint32_t>(blockIdx.x), static_cast<__hip_uint32_t>(blockIdx.y), - static_cast<__hip_uint32_t>(blockIdx.z))); + return (dim3(static_cast(blockIdx.x), static_cast(blockIdx.y), + static_cast(blockIdx.z))); } __CG_STATIC_QUALIFIER__ dim3 thread_index() { - return (dim3(static_cast<__hip_uint32_t>(threadIdx.x), static_cast<__hip_uint32_t>(threadIdx.y), - static_cast<__hip_uint32_t>(threadIdx.z))); + return (dim3(static_cast(threadIdx.x), static_cast(threadIdx.y), + static_cast(threadIdx.z))); } -__CG_STATIC_QUALIFIER__ __hip_uint32_t size() { - return (static_cast<__hip_uint32_t>(blockDim.x * blockDim.y * blockDim.z)); +__CG_STATIC_QUALIFIER__ uint32_t size() { + return (static_cast(blockDim.x * blockDim.y * blockDim.z)); } -__CG_STATIC_QUALIFIER__ __hip_uint32_t thread_rank() { - return (static_cast<__hip_uint32_t>((threadIdx.z * blockDim.y * blockDim.x) + +__CG_STATIC_QUALIFIER__ uint32_t thread_rank() { + return (static_cast((threadIdx.z * blockDim.y * blockDim.x) + (threadIdx.y * blockDim.x) + (threadIdx.x))); } @@ -196,8 +196,8 @@ __CG_STATIC_QUALIFIER__ bool is_valid() { __CG_STATIC_QUALIFIER__ void sync() { __syncthreads(); } __CG_STATIC_QUALIFIER__ dim3 block_dim() { - return (dim3(static_cast<__hip_uint32_t>(blockDim.x), static_cast<__hip_uint32_t>(blockDim.y), - static_cast<__hip_uint32_t>(blockDim.z))); + return (dim3(static_cast(blockDim.x), static_cast(blockDim.y), + static_cast(blockDim.z))); } } // namespace workgroup diff --git a/projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake b/projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake index ef9e559d1f..680e4da429 100644 --- a/projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake +++ b/projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake @@ -60,6 +60,13 @@ function(get_hiprtc_macros HIPRTC_DEFINES) #define HIP_INCLUDE_HIP_MATH_FUNCTIONS_H\n\ #define HIP_INCLUDE_HIP_HIP_VECTOR_TYPES_H\n\ #if !__HIP_NO_STD_DEFS__\n\ +#if defined(_WIN32)\n\ +typedef unsigned long long uint64_t;\n\ +typedef signed long long int64_t;\n\ +#else\n\ +typedef unsigned long uint64_t;\n\ +typedef signed long int64_t;\n\ +#endif\n\ #if defined(__HIPRTC_PTRDIFF_T_IS_LONG_LONG__) && __HIPRTC_PTRDIFF_T_IS_LONG_LONG__==1\n\ typedef long long ptrdiff_t;\n\ #else\n\ @@ -67,6 +74,8 @@ typedef __PTRDIFF_TYPE__ ptrdiff_t;\n\ #endif\n\ typedef long clock_t;\n\ namespace std {\n\ +using ::uint64_t;\n\ +using ::int64_t;\n\ using ::ptrdiff_t;\n\ using ::clock_t;\n\ }\n\