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

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

Browse Source 
This reverts commit 73558e3363.

Reason for revert: <INSERT REASONING HERE>

Change-Id: I53322718dadde2c98f96140b8e260ec7ee9ef721
This commit is contained in:
Vladislav Sytchenko
2020-10-05 12:26:18 -04:00
parent 23a9cb15eb
commit 374ead1d19
111 changed files with 750 additions and 6797 deletions
Download Patch File Download Diff File Expand all files Collapse all files
include/hip/hcc_detail/device_functions.h
+2 -136
View File
@@ -365,25 +365,6 @@ 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), "");
@@ -397,22 +378,8 @@ 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), "");
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__
__device__
inline
int __shfl_up(int var, unsigned int lane_delta, int width = warpSize) {
int self = __lane_id();
@@ -468,28 +435,6 @@ 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)
@@ -504,20 +449,6 @@ 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) {
@@ -576,26 +507,6 @@ 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), "");
@@ -607,19 +518,6 @@ 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
@@ -679,26 +577,6 @@ 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), "");
@@ -710,19 +588,7 @@ 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,22 +487,6 @@ 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
+5 -8
View File
@@ -345,16 +345,13 @@ typedef struct hipLaunchParams_t {
hipStream_t stream; ///< Stream identifier
} hipLaunchParams;
#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);
// 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);
#endif
// Doxygen end group GlobalDefs
/** @} */
include/hip/hip_cooperative_groups.h
+3 -6
View File
@@ -28,17 +28,14 @@ THE SOFTWARE.
*/
#ifndef HIP_INCLUDE_HIP_HIP_COOPERATIVE_GROUP_H
#define HIP_INCLUDE_HIP_HIP_COOPERATIVE_GROUP_H
#include <hip/hip_version.h>
#include <hip/hip_common.h>
#define HIP_INCLUDE_HIP_HIP_VECTOR_TYPES_H
#if defined(__HIP_PLATFORM_HCC__) && !defined(__HIP_PLATFORM_NVCC__)
#if __cplusplus && defined(__clang__) && defined(__HIP__)
#if __cplusplus
#include <hip/hcc_detail/hip_cooperative_groups.h>
#endif
#elif defined(__HIP_PLATFORM_NVCC__) && !defined(__HIP_PLATFORM_HCC__)
#include <hip/nvcc_detail/hip_cooperative_groups.h>
#include <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,7 +32,6 @@ 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
@@ -1,12 +0,0 @@
#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,7 +26,6 @@ THE SOFTWARE.
#include <cuda_runtime_api.h>
#include <cuda.h>
#include <cuda_profiler_api.h>
#include <cuda_fp16.h>
#ifdef __cplusplus
extern "C" {