AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

SWDEV-431175 - Remove datatypes from HIPRTC

Просмотр исходного кода 
Defining int64_t, uint64_t, int32_t, uint32_t in HIPRTC
seem to result in conflicts with some apps as they use
their own definitions for these types. NVRTC also doesn't
define these. Hence remove them to match the behavior.

Change-Id: I77ef70e846950698cb00375f5d0501b907f01fe3


[ROCm/clr commit: 629e279f72]
Этот коммит содержится в:
Satyanvesh Dittakavi
2023-11-07 18:44:32 +00:00
коммит произвёл Rakesh Roy
родитель cc6e0c2a32
Коммит 6035b47801
8 изменённых файлов: 168 добавлений и 185 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h
+41 -41
Просмотреть файл
@@ -59,7 +59,7 @@ __device__ static inline int __clz(int input) {
}
__device__ static inline int __clzll(long long int input) {
return __ockl_clz_u64((uint64_t)input);
return __ockl_clz_u64((__hip_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 int32_t __fns64(uint64_t mask, uint32_t base, int32_t offset) {
uint64_t temp_mask = mask;
int32_t temp_offset = offset;
__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;
if (offset == 0) {
temp_mask &= (1 << base);
@@ -98,10 +98,10 @@ __device__ static int32_t __fns64(uint64_t mask, uint32_t base, int32_t offset)
temp_mask = temp_mask & ((~0ULL) << base);
if (__builtin_popcountll(temp_mask) < temp_offset)
return -1;
int32_t total = 0;
__hip_int32_t total = 0;
for (int i = 0x20; i > 0; i >>= 1) {
uint64_t temp_mask_lo = temp_mask & ((1ULL << i) - 1);
int32_t pcnt = __builtin_popcountll(temp_mask_lo);
__hip_uint64_t temp_mask_lo = temp_mask & ((1ULL << i) - 1);
__hip_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 int32_t __fns64(uint64_t mask, uint32_t base, int32_t offset)
return total;
}
__device__ static int32_t __fns32(uint64_t mask, uint32_t base, int32_t offset) {
uint64_t temp_mask = mask;
int32_t temp_offset = offset;
__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;
if (offset == 0) {
temp_mask &= (1 << base);
temp_offset = 1;
@@ -132,10 +132,10 @@ __device__ static int32_t __fns32(uint64_t mask, uint32_t base, int32_t offset)
temp_mask = temp_mask & ((~0ULL) << base);
if (__builtin_popcountll(temp_mask) < temp_offset)
return -1;
int32_t total = 0;
__hip_int32_t total = 0;
for (int i = 0x20; i > 0; i >>= 1) {
uint64_t temp_mask_lo = temp_mask & ((1ULL << i) - 1);
int32_t pcnt = __builtin_popcountll(temp_mask_lo);
__hip_uint64_t temp_mask_lo = temp_mask & ((1ULL << i) - 1);
__hip_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(uint64_t input) {
__device__ static inline unsigned int __lastbit_u32_u64(__hip_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) {
uint32_t offset = src1 & 31;
uint32_t width = src2 & 31;
__hip_uint32_t offset = src1 & 31;
__hip_uint32_t width = src2 & 31;
return width == 0 ? 0 : (src0 << (32 - offset - width)) >> (32 - width);
}
__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;
__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;
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) {
uint32_t offset = src2 & 31;
uint32_t width = src3 & 31;
uint32_t mask = (1 << width) - 1;
__hip_uint32_t offset = src2 & 31;
__hip_uint32_t width = src3 & 31;
__hip_uint32_t mask = (1 << width) - 1;
return ((src0 & ~(mask << offset)) | ((src1 & mask) << offset));
}
__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;
__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;
return ((src0 & ~(mask << offset)) | ((src1 & mask) << offset));
}
__device__ inline unsigned int __funnelshift_l(unsigned int lo, unsigned int hi, unsigned int shift)
{
uint32_t mask_shift = shift & 31;
__hip_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)
{
uint32_t min_shift = shift >= 32 ? 32 : shift;
__hip_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(uint64_t), "");
static_assert(sizeof(double) == sizeof(__hip_uint64_t), "");
uint64_t tmp0 = (static_cast<uint64_t>(hi) << 32ull) | static_cast<uint32_t>(lo);
__hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(hi) << 32ull) | static_cast<__hip_uint32_t>(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
uint64_t __lanemask_gt()
__hip_uint64_t __lanemask_gt()
{
uint32_t lane = __ockl_lane_u32();
__hip_uint32_t lane = __ockl_lane_u32();
if (lane == 63)
return 0;
uint64_t ballot = __ballot64(1);
uint64_t mask = (~((uint64_t)0)) << (lane + 1);
__hip_uint64_t ballot = __ballot64(1);
__hip_uint64_t mask = (~((__hip_uint64_t)0)) << (lane + 1);
return mask & ballot;
}
__device__
inline
uint64_t __lanemask_lt()
__hip_uint64_t __lanemask_lt()
{
uint32_t lane = __ockl_lane_u32();
int64_t ballot = __ballot64(1);
uint64_t mask = ((uint64_t)1 << lane) - (uint64_t)1;
__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;
return mask & ballot;
}
__device__
inline
uint64_t __lanemask_eq()
__hip_uint64_t __lanemask_eq()
{
uint32_t lane = __ockl_lane_u32();
int64_t mask = ((uint64_t)1 << lane);
__hip_uint32_t lane = __ockl_lane_u32();
__hip_int64_t mask = ((__hip_uint64_t)1 << lane);
return mask;
}
projects/clr/hipamd/include/hip/amd_detail/amd_hip_bfloat16.h
+4 -4
Просмотреть файл
@@ -81,9 +81,9 @@ struct hip_bfloat16
{
union
{
uint32_t int32;
__hip_uint32_t int32;
float fp32;
} u = {uint32_t(data) << 16};
} u = {__hip_uint32_t(data) << 16};
return u.fp32;
}
@@ -113,7 +113,7 @@ private:
union
{
float fp32;
uint32_t int32;
__hip_uint32_t int32;
} u = {f};
if(~u.int32 & 0x7f800000)
{
@@ -156,7 +156,7 @@ private:
union
{
float fp32;
uint32_t int32;
__hip_uint32_t int32;
} u = {f};
return __hip_uint16_t(u.int32 >> 16) | (!(~u.int32 & 0x7f800000) && (u.int32 & 0xffff));
}
projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h
+19 -19
Просмотреть файл
@@ -62,9 +62,9 @@ namespace cooperative_groups {
*/
class thread_group {
protected:
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,
__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,
// 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, uint32_t size = static_cast<uint64_t>(0),
uint64_t mask = static_cast<uint64_t>(0)) {
__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)) {
_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__ uint32_t size() const { return _size; }
__CG_QUALIFIER__ __hip_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__ uint32_t thread_rank() const;
__CG_QUALIFIER__ __hip_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(uint32_t size)
explicit __CG_QUALIFIER__ multi_grid_group(__hip_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__ uint32_t num_grids() { return internal::multi_grid::num_grids(); }
__CG_QUALIFIER__ __hip_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__ 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__ __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__ 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(uint32_t size) : thread_group(internal::cg_grid, size) {}
explicit __CG_QUALIFIER__ grid_group(__hip_uint32_t size) : thread_group(internal::cg_grid, size) {}
public:
__CG_QUALIFIER__ uint32_t thread_rank() const { return internal::grid::thread_rank(); }
__CG_QUALIFIER__ __hip_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(uint32_t size)
explicit __CG_QUALIFIER__ thread_block(__hip_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__ uint32_t thread_rank() { return internal::workgroup::thread_rank(); }
__CG_STATIC_QUALIFIER__ uint32_t size() { return internal::workgroup::size(); }
__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__ 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__ uint32_t thread_group::thread_rank() const {
__CG_QUALIFIER__ __hip_uint32_t thread_group::thread_rank() const {
switch (this->_type) {
case internal::cg_multi_grid: {
return (static_cast<const multi_grid_group*>(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 <class CGTy> __CG_QUALIFIER__ uint32_t group_size(CGTy const& g) { return g.size(); }
template <class CGTy> __CG_QUALIFIER__ __hip_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 <class CGTy> __CG_QUALIFIER__ uint32_t thread_rank(CGTy const& g) {
template <class CGTy> __CG_QUALIFIER__ __hip_uint32_t thread_rank(CGTy const& g) {
return g.thread_rank();
}
/**
projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime.h
+15 -23
Просмотреть файл
@@ -87,15 +87,6 @@ size_t amd_dbgapi_get_build_id();
#include <math.h>
#include <stdint.h>
#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__
@@ -249,34 +240,35 @@ 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__
std::uint32_t operator()(std::uint32_t x) const noexcept { return __ockl_get_group_id(x); }
__hip_uint32_t operator()(__hip_uint32_t x) const noexcept { return __ockl_get_group_id(x); }
};
struct __HIP_BlockDim {
__device__
std::uint32_t operator()(std::uint32_t x) const noexcept {
__hip_uint32_t operator()(__hip_uint32_t x) const noexcept {
return __ockl_get_local_size(x);
}
};
struct __HIP_GridDim {
__device__
std::uint32_t operator()(std::uint32_t x) const noexcept {
__hip_uint32_t operator()(__hip_uint32_t x) const noexcept {
return __ockl_get_num_groups(x);
}
};
struct __HIP_ThreadIdx {
__device__
std::uint32_t operator()(std::uint32_t x) const noexcept {
__hip_uint32_t operator()(__hip_uint32_t x) const noexcept {
return __ockl_get_local_id(x);
}
};
#if defined(__HIPCC_RTC__)
typedef struct dim3 {
uint32_t x; ///< x
uint32_t y; ///< y
uint32_t z; ///< z
__hip_uint32_t x; ///< x
__hip_uint32_t y; ///< y
__hip_uint32_t z; ///< z
#ifdef __cplusplus
constexpr __device__ dim3(uint32_t _x = 1, uint32_t _y = 1, uint32_t _z = 1) : x(_x), y(_y), z(_z){};
constexpr __device__ dim3(__hip_uint32_t _x = 1, __hip_uint32_t _y = 1, __hip_uint32_t _z = 1)
: x(_x), y(_y), z(_z){};
#endif
} dim3;
#endif // !defined(__HIPCC_RTC__)
@@ -325,37 +317,37 @@ constexpr typename __HIP_Coordinates<F>::__Z __HIP_Coordinates<F>::z;
inline
__device__
std::uint32_t operator*(__HIP_Coordinates<__HIP_GridDim>::__X,
__hip_uint32_t operator*(__HIP_Coordinates<__HIP_GridDim>::__X,
__HIP_Coordinates<__HIP_BlockDim>::__X) noexcept {
return __ockl_get_global_size(0);
}
inline
__device__
std::uint32_t operator*(__HIP_Coordinates<__HIP_BlockDim>::__X,
__hip_uint32_t operator*(__HIP_Coordinates<__HIP_BlockDim>::__X,
__HIP_Coordinates<__HIP_GridDim>::__X) noexcept {
return __ockl_get_global_size(0);
}
inline
__device__
std::uint32_t operator*(__HIP_Coordinates<__HIP_GridDim>::__Y,
__hip_uint32_t operator*(__HIP_Coordinates<__HIP_GridDim>::__Y,
__HIP_Coordinates<__HIP_BlockDim>::__Y) noexcept {
return __ockl_get_global_size(1);
}
inline
__device__
std::uint32_t operator*(__HIP_Coordinates<__HIP_BlockDim>::__Y,
__hip_uint32_t operator*(__HIP_Coordinates<__HIP_BlockDim>::__Y,
__HIP_Coordinates<__HIP_GridDim>::__Y) noexcept {
return __ockl_get_global_size(1);
}
inline
__device__
std::uint32_t operator*(__HIP_Coordinates<__HIP_GridDim>::__Z,
__hip_uint32_t operator*(__HIP_Coordinates<__HIP_GridDim>::__Z,
__HIP_Coordinates<__HIP_BlockDim>::__Z) noexcept {
return __ockl_get_global_size(2);
}
inline
__device__
std::uint32_t operator*(__HIP_Coordinates<__HIP_BlockDim>::__Z,
__hip_uint32_t operator*(__HIP_Coordinates<__HIP_BlockDim>::__Z,
__HIP_Coordinates<__HIP_GridDim>::__Z) noexcept {
return __ockl_get_global_size(2);
}
projects/clr/hipamd/include/hip/amd_detail/amd_warp_functions.h
+40 -40
Просмотреть файл
@@ -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(uint64_t), "");
static_assert(sizeof(double) == sizeof(__hip_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);
uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
__hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(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(uint64_t), "");
static_assert(sizeof(long) == sizeof(__hip_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);
uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
__hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(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(uint64_t), "");
static_assert(sizeof(unsigned long) == sizeof(__hip_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);
uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
__hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(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(uint64_t), "");
static_assert(sizeof(long long) == sizeof(__hip_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);
uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
__hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(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(uint64_t), "");
static_assert(sizeof(unsigned long long) == sizeof(__hip_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);
uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
__hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(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(uint64_t), "");
static_assert(sizeof(double) == sizeof(__hip_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);
uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
__hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(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(uint64_t), "");
static_assert(sizeof(long) == sizeof(__hip_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);
uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
__hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(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(uint64_t), "");
static_assert(sizeof(unsigned long) == sizeof(__hip_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);
uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
__hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(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(uint64_t), "");
static_assert(sizeof(long long) == sizeof(__hip_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);
uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
__hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(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(uint64_t), "");
static_assert(sizeof(unsigned long long) == sizeof(__hip_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);
uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
__hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(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(uint64_t), "");
static_assert(sizeof(double) == sizeof(__hip_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);
uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
__hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(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(uint64_t), "");
static_assert(sizeof(long) == sizeof(__hip_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);
uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
__hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(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(uint64_t), "");
static_assert(sizeof(unsigned long) == sizeof(__hip_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);
uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
__hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(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(uint64_t), "");
static_assert(sizeof(long long) == sizeof(__hip_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);
uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
__hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(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(uint64_t), "");
static_assert(sizeof(unsigned long long) == sizeof(__hip_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);
uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
__hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(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(uint64_t), "");
static_assert(sizeof(double) == sizeof(__hip_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);
uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
__hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(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(uint64_t), "");
static_assert(sizeof(long) == sizeof(__hip_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);
uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
__hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(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(uint64_t), "");
static_assert(sizeof(unsigned long) == sizeof(__hip_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);
uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
__hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(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(uint64_t), "");
static_assert(sizeof(long long) == sizeof(__hip_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);
uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
__hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(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(uint64_t), "");
static_assert(sizeof(unsigned long long) == sizeof(__hip_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);
uint64_t tmp0 = (static_cast<uint64_t>(tmp[1]) << 32ull) | static_cast<uint32_t>(tmp[0]);
__hip_uint64_t tmp0 = (static_cast<__hip_uint64_t>(tmp[1]) << 32ull) | static_cast<__hip_uint32_t>(tmp[0]);
unsigned long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0));
return tmp1;
}
projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h
+25 -25
Просмотреть файл
@@ -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)) uint64_t __ockl_clz_u64(uint64_t);
extern "C" __device__ __attribute__((const)) __hip_uint64_t __ockl_clz_u64(__hip_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(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__((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__((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)) uint32_t __ockl_lane_u32();
extern "C" __device__ __attribute__((const)) __hip_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__ 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);
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);
// Introduce local address space
#define __local __attribute__((address_space(3)))
projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h
+24 -24
Просмотреть файл
@@ -112,16 +112,16 @@ typedef enum {
*/
namespace multi_grid {
__CG_STATIC_QUALIFIER__ uint32_t num_grids() {
return static_cast<uint32_t>(__ockl_multi_grid_num_grids()); }
__CG_STATIC_QUALIFIER__ __hip_uint32_t num_grids() {
return static_cast<__hip_uint32_t>(__ockl_multi_grid_num_grids()); }
__CG_STATIC_QUALIFIER__ uint32_t grid_rank() {
return static_cast<uint32_t>(__ockl_multi_grid_grid_rank()); }
__CG_STATIC_QUALIFIER__ __hip_uint32_t grid_rank() {
return static_cast<__hip_uint32_t>(__ockl_multi_grid_grid_rank()); }
__CG_STATIC_QUALIFIER__ uint32_t size() { return static_cast<uint32_t>(__ockl_multi_grid_size()); }
__CG_STATIC_QUALIFIER__ __hip_uint32_t size() { return static_cast<__hip_uint32_t>(__ockl_multi_grid_size()); }
__CG_STATIC_QUALIFIER__ uint32_t thread_rank() {
return static_cast<uint32_t>(__ockl_multi_grid_thread_rank()); }
__CG_STATIC_QUALIFIER__ __hip_uint32_t thread_rank() {
return static_cast<__hip_uint32_t>(__ockl_multi_grid_thread_rank()); }
__CG_STATIC_QUALIFIER__ bool is_valid() { return static_cast<bool>(__ockl_multi_grid_is_valid()); }
@@ -135,23 +135,23 @@ __CG_STATIC_QUALIFIER__ void sync() { __ockl_multi_grid_sync(); }
*/
namespace grid {
__CG_STATIC_QUALIFIER__ uint32_t size() {
return static_cast<uint32_t>((blockDim.z * gridDim.z) * (blockDim.y * gridDim.y) *
__CG_STATIC_QUALIFIER__ __hip_uint32_t size() {
return static_cast<__hip_uint32_t>((blockDim.z * gridDim.z) * (blockDim.y * gridDim.y) *
(blockDim.x * gridDim.x));
}
__CG_STATIC_QUALIFIER__ uint32_t thread_rank() {
__CG_STATIC_QUALIFIER__ __hip_uint32_t thread_rank() {
// Compute global id of the workgroup to which the current thread belongs to
uint32_t blkIdx = static_cast<uint32_t>((blockIdx.z * gridDim.y * gridDim.x) +
__hip_uint32_t blkIdx = static_cast<__hip_uint32_t>((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
uint32_t num_threads_till_current_workgroup =
static_cast<uint32_t>(blkIdx * (blockDim.x * blockDim.y * blockDim.z));
__hip_uint32_t num_threads_till_current_workgroup =
static_cast<__hip_uint32_t>(blkIdx * (blockDim.x * blockDim.y * blockDim.z));
// Compute thread local rank within current workgroup
uint32_t local_thread_rank = static_cast<uint32_t>((threadIdx.z * blockDim.y * blockDim.x) +
__hip_uint32_t local_thread_rank = static_cast<__hip_uint32_t>((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<uint32_t>(blockIdx.x), static_cast<uint32_t>(blockIdx.y),
static_cast<uint32_t>(blockIdx.z)));
return (dim3(static_cast<__hip_uint32_t>(blockIdx.x), static_cast<__hip_uint32_t>(blockIdx.y),
static_cast<__hip_uint32_t>(blockIdx.z)));
}
__CG_STATIC_QUALIFIER__ dim3 thread_index() {
return (dim3(static_cast<uint32_t>(threadIdx.x), static_cast<uint32_t>(threadIdx.y),
static_cast<uint32_t>(threadIdx.z)));
return (dim3(static_cast<__hip_uint32_t>(threadIdx.x), static_cast<__hip_uint32_t>(threadIdx.y),
static_cast<__hip_uint32_t>(threadIdx.z)));
}
__CG_STATIC_QUALIFIER__ uint32_t size() {
return (static_cast<uint32_t>(blockDim.x * blockDim.y * blockDim.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 thread_rank() {
return (static_cast<uint32_t>((threadIdx.z * blockDim.y * blockDim.x) +
__CG_STATIC_QUALIFIER__ __hip_uint32_t thread_rank() {
return (static_cast<__hip_uint32_t>((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<uint32_t>(blockDim.x), static_cast<uint32_t>(blockDim.y),
static_cast<uint32_t>(blockDim.z)));
return (dim3(static_cast<__hip_uint32_t>(blockDim.x), static_cast<__hip_uint32_t>(blockDim.y),
static_cast<__hip_uint32_t>(blockDim.z)));
}
} // namespace workgroup
projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake
-9
Просмотреть файл
@@ -60,13 +60,6 @@ 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\
@@ -74,8 +67,6 @@ 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\