AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Revert "Revert "Merge branch 'amd-master-next' into amd-npi-next""

Browse Source 
This reverts commit 374ead1d19.

Reason for revert: <INSERT REASONING HERE>

Change-Id: I92ceb171e31026ed1864704cef2fc1497b883ef9
This commit is contained in:
Vladislav Sytchenko
2020-10-05 13:20:58 -04:00
parent 374ead1d19
commit ad2d55c144
111 changed files with 6800 additions and 753 deletions
Download Patch File Download Diff File Expand all files Collapse all files
include/hip/hcc_detail/device_functions.h
+136 -2
View File
@@ -365,6 +365,25 @@ long __shfl(long var, int src_lane, int width = warpSize)
}
__device__
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), "");
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]);
unsigned long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0));
return tmp1;
#else
static_assert(sizeof(unsigned long) == sizeof(unsigned int), "");
return static_cast<unsigned long>(__shfl(static_cast<unsigned int>(var), src_lane, width));
#endif
}
__device__
inline
long long __shfl(long long var, int src_lane, int width = warpSize)
{
static_assert(sizeof(long long) == 2 * sizeof(int), "");
@@ -378,8 +397,22 @@ long long __shfl(long long var, int src_lane, int width = warpSize)
long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0));
return tmp1;
}
__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), "");
__device__
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]);
unsigned long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0));
return tmp1;
}
__device__
inline
int __shfl_up(int var, unsigned int lane_delta, int width = warpSize) {
int self = __lane_id();
@@ -435,6 +468,28 @@ long __shfl_up(long var, unsigned int lane_delta, int width = warpSize)
return static_cast<long>(__shfl_up(static_cast<int>(var), lane_delta, width));
#endif
}
__device__
inline
unsigned long __shfl_up(unsigned long var, unsigned int lane_delta, int width = warpSize)
{
#ifndef _MSC_VER
static_assert(sizeof(unsigned long) == 2 * sizeof(unsigned int), "");
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);
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
static_assert(sizeof(unsigned long) == sizeof(unsigned int), "");
return static_cast<unsigned long>(__shfl_up(static_cast<unsigned int>(var), lane_delta, width));
#endif
}
__device__
inline
long long __shfl_up(long long var, unsigned int lane_delta, int width = warpSize)
@@ -449,6 +504,20 @@ long long __shfl_up(long long var, unsigned int lane_delta, int width = warpSize
return tmp1;
}
__device__
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), "");
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]);
unsigned long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0));
return tmp1;
}
__device__
inline
int __shfl_down(int var, unsigned int lane_delta, int width = warpSize) {
@@ -507,6 +576,26 @@ long __shfl_down(long var, unsigned int lane_delta, int width = warpSize)
}
__device__
inline
unsigned long __shfl_down(unsigned long var, unsigned int lane_delta, int width = warpSize)
{
#ifndef _MSC_VER
static_assert(sizeof(unsigned long) == 2 * sizeof(unsigned int), "");
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);
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
static_assert(sizeof(unsigned long) == sizeof(unsigned int), "");
return static_cast<unsigned long>(__shfl_down(static_cast<unsigned int>(var), lane_delta, width));
#endif
}
__device__
inline
long long __shfl_down(long long var, unsigned int lane_delta, int width = warpSize)
{
static_assert(sizeof(long long) == 2 * sizeof(int), "");
@@ -518,6 +607,19 @@ long long __shfl_down(long long var, unsigned int lane_delta, int width = warpSi
long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0));
return tmp1;
}
__device__
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), "");
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]);
unsigned long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0));
return tmp1;
}
__device__
inline
@@ -577,6 +679,26 @@ long __shfl_xor(long var, int lane_mask, int width = warpSize)
}
__device__
inline
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), "");
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]);
unsigned long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0));
return tmp1;
#else
static_assert(sizeof(unsigned long) == sizeof(unsigned int), "");
return static_cast<unsigned long>(__shfl_xor(static_cast<unsigned int>(var), lane_mask, width));
#endif
}
__device__
inline
long long __shfl_xor(long long var, int lane_mask, int width = warpSize)
{
static_assert(sizeof(long long) == 2 * sizeof(int), "");
@@ -588,7 +710,19 @@ long long __shfl_xor(long long var, int lane_mask, int width = warpSize)
long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0));
return tmp1;
}
__device__
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), "");
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]);
unsigned long long tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0));
return tmp1;
}
#define MASK1 0x00ff00ff
#define MASK2 0xff00ff00
include/hip/hcc_detail/hip_runtime.h
+16
View File
@@ -487,6 +487,22 @@ struct __HIP_Coordinates {
#endif
};
template <typename F>
#if !defined(_MSC_VER)
__attribute__((weak))
#endif
constexpr typename __HIP_Coordinates<F>::X __HIP_Coordinates<F>::x;
template <typename F>
#if !defined(_MSC_VER)
__attribute__((weak))
#endif
constexpr typename __HIP_Coordinates<F>::Y __HIP_Coordinates<F>::y;
template <typename F>
#if !defined(_MSC_VER)
__attribute__((weak))
#endif
constexpr typename __HIP_Coordinates<F>::Z __HIP_Coordinates<F>::z;
extern "C" __device__ __attribute__((const)) size_t __ockl_get_global_size(uint);
inline
__device__
include/hip/hcc_detail/hip_runtime_api.h
+8 -5
View File
@@ -345,13 +345,16 @@ typedef struct hipLaunchParams_t {
hipStream_t stream; ///< Stream identifier
} hipLaunchParams;
// Pre-Compiled header for online compilation
#ifdef ENABLE_HIP_PCH
extern const char* __hip_pch;
extern unsigned __hip_pch_size;
void __hipGetPCH(const char** pch, unsigned int*size);
#if __HIP_HAS_GET_PCH
/**
* Internal use only. This API may change in the future
* Pre-Compiled header for online compilation
*
*/
void __hipGetPCH(const char** pch, unsigned int*size);
#endif
// Doxygen end group GlobalDefs
/** @} */
include/hip/hip_cooperative_groups.h
+6 -3
View File
@@ -28,14 +28,17 @@ THE SOFTWARE.
*/
#ifndef HIP_INCLUDE_HIP_HIP_COOPERATIVE_GROUP_H
#define HIP_INCLUDE_HIP_HIP_VECTOR_TYPES_H
#define HIP_INCLUDE_HIP_HIP_COOPERATIVE_GROUP_H
#include <hip/hip_version.h>
#include <hip/hip_common.h>
#if defined(__HIP_PLATFORM_HCC__) && !defined(__HIP_PLATFORM_NVCC__)
#if __cplusplus
#if __cplusplus && defined(__clang__) && defined(__HIP__)
#include <hip/hcc_detail/hip_cooperative_groups.h>
#endif
#elif defined(__HIP_PLATFORM_NVCC__) && !defined(__HIP_PLATFORM_HCC__)
#include <cooperative_groups.h>
#include <hip/nvcc_detail/hip_cooperative_groups.h>
#else
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
#endif
include/hip/hip_runtime_api.h
+1
View File
@@ -32,6 +32,7 @@ THE SOFTWARE.
#include <string.h> // for getDeviceProp
#include <hip/hip_version.h>
#include <hip/hip_common.h>
enum {
include/hip/nvcc_detail/hip_cooperative_groups.h
+12
View File
@@ -0,0 +1,12 @@
#ifndef HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_COOPERATIVE_GROUPS_H
#define HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_COOPERATIVE_GROUPS_H
// Include CUDA headers
#include <cuda_runtime.h>
#include <cooperative_groups.h>
// Include HIP wrapper headers around CUDA
#include <hip/hip_runtime.h>
#include <hip/hip_runtime_api.h>
#endif // HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_COOPERATIVE_GROUPS_H
include/hip/nvcc_detail/hip_runtime.h
+2 -2
View File
@@ -104,13 +104,13 @@ typedef int hipLaunchParm;
#define HIP_DYNAMIC_SHARED_ATTRIBUTE
#ifdef __HIP_DEVICE_COMPILE__
#define abort() \
#define abort_() \
{ asm("trap;"); }
#undef assert
#define assert(COND) \
{ \
if (!COND) { \
abort(); \
abort_(); \
} \
}
#endif
include/hip/nvcc_detail/hip_runtime_api.h
+1
View File
@@ -26,6 +26,7 @@ THE SOFTWARE.
#include <cuda_runtime_api.h>
#include <cuda.h>
#include <cuda_profiler_api.h>
#include <cuda_fp16.h>
#ifdef __cplusplus
extern "C" {