AI/rocm-systems
2
0
Forkuj 0
Kod Zgłoszenia Oczekujące zmiany Actions Packages Projekty Wydania Wiki Aktywność

Revert "SWDEV-431175 - Remove datatypes from HIPRTC"

Przeglądaj źródła 
This reverts commit 6035b47801.

Reason for revert: This is considered a breaking change and requires
multiple apps to change their behavior. This will be reintroduced in later releases.

Change-Id: I0354ce4e0f5e6c402499a7a8c2aaf43bf5b1bfc7


[ROCm/clr commit: 270c6d48f1]
This commit is contained in:
Satyanvesh Dittakavi
2023-12-01 07:30:11 +00:00
zatwierdzone przez Maneesh Gupta
rodzic 0d74d8b187
commit 85e9925335
8 zmienionych plików z 185 dodań i 168 usunięć
Ściągnij plik aktualizacji Ściągnij plik porównania Expand all files Collapse all files
projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h
+41 -41
Wyświetl plik
@@ -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<uint64_t>(hi) << 32ull) | static_cast<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
__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;
}
projects/clr/hipamd/include/hip/amd_detail/amd_hip_bfloat16.h
+4 -4
Wyświetl plik
@@ -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));
}
projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h
+19 -19
Wyświetl plik
@@ -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<uint64_t>(0),
uint64_t mask = static_cast<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__ __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<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__ __hip_uint32_t group_size(CGTy const& g) { return g.size(); }
template <class CGTy> __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 <class CGTy> __CG_QUALIFIER__ __hip_uint32_t thread_rank(CGTy const& g) {
template <class CGTy> __CG_QUALIFIER__ uint32_t thread_rank(CGTy const& g) {
return g.thread_rank();
}
/**
projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime.h
+23 -15
Wyświetl plik
@@ -87,6 +87,15 @@ 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__
@@ -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<F>::__Z __HIP_Coordinates<F>::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);
}
projects/clr/hipamd/include/hip/amd_detail/amd_warp_functions.h
+40 -40
Wyświetl plik
@@ -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<uint64_t>(tmp[1]) << 32ull) | static_cast<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(__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<uint64_t>(tmp[1]) << 32ull) | static_cast<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(__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<uint64_t>(tmp[1]) << 32ull) | static_cast<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(__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<uint64_t>(tmp[1]) << 32ull) | static_cast<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(__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<uint64_t>(tmp[1]) << 32ull) | static_cast<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(__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<uint64_t>(tmp[1]) << 32ull) | static_cast<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(__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<uint64_t>(tmp[1]) << 32ull) | static_cast<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(__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<uint64_t>(tmp[1]) << 32ull) | static_cast<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(__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<uint64_t>(tmp[1]) << 32ull) | static_cast<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(__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<uint64_t>(tmp[1]) << 32ull) | static_cast<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(__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<uint64_t>(tmp[1]) << 32ull) | static_cast<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(__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<uint64_t>(tmp[1]) << 32ull) | static_cast<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(__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<uint64_t>(tmp[1]) << 32ull) | static_cast<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(__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<uint64_t>(tmp[1]) << 32ull) | static_cast<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(__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<uint64_t>(tmp[1]) << 32ull) | static_cast<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(__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<uint64_t>(tmp[1]) << 32ull) | static_cast<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(__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<uint64_t>(tmp[1]) << 32ull) | static_cast<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(__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<uint64_t>(tmp[1]) << 32ull) | static_cast<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(__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<uint64_t>(tmp[1]) << 32ull) | static_cast<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(__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<uint64_t>(tmp[1]) << 32ull) | static_cast<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
Wyświetl plik
@@ -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)))
projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h
+24 -24
Wyświetl plik
@@ -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<uint32_t>(__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<uint32_t>(__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<uint32_t>(__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<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__ __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<uint32_t>((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<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
__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<uint32_t>(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<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<__hip_uint32_t>(blockIdx.x), static_cast<__hip_uint32_t>(blockIdx.y),
static_cast<__hip_uint32_t>(blockIdx.z)));
return (dim3(static_cast<uint32_t>(blockIdx.x), static_cast<uint32_t>(blockIdx.y),
static_cast<uint32_t>(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<uint32_t>(threadIdx.x), static_cast<uint32_t>(threadIdx.y),
static_cast<uint32_t>(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<uint32_t>(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<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<__hip_uint32_t>(blockDim.x), static_cast<__hip_uint32_t>(blockDim.y),
static_cast<__hip_uint32_t>(blockDim.z)));
return (dim3(static_cast<uint32_t>(blockDim.x), static_cast<uint32_t>(blockDim.y),
static_cast<uint32_t>(blockDim.z)));
}
} // namespace workgroup
projects/clr/hipamd/src/hiprtc/cmake/HIPRTC.cmake
+9
Wyświetl plik
@@ -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\