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

Merge branch 'master' into support-malloc

Przeglądaj źródła
This commit is contained in:
Maneesh Gupta
2018-09-17 10:17:25 +05:30
zatwierdzone przez GitHub
rodzic f06894e2f0 48cc4f0d9e
commit 66f863d1f3
61 zmienionych plików z 5094 dodań i 2387 usunięć
Ściągnij plik aktualizacji Ściągnij plik porównania Expand all files Collapse all files
include/hip/hcc_detail/device_functions.h
+70 -63
Wyświetl plik
@@ -32,9 +32,6 @@ THE SOFTWARE.
#include <hip/hcc_detail/llvm_intrinsics.h>
#include <stddef.h>
typedef unsigned long ulong;
typedef unsigned int uint;
/*
Integer Intrinsics
*/
@@ -47,78 +44,68 @@ __device__ static inline unsigned int __popcll(unsigned long long int input) {
return __builtin_popcountl(input);
}
__device__ static inline unsigned int __clz(unsigned int input) {
#ifdef NVCC_COMPAT
return input == 0 ? 32 : __builtin_clz(input);
#else
return input == 0 ? -1 : __builtin_clz(input);
#endif
__device__ static inline int __clz(int input) {
return __ockl_clz_u32((uint)input);
}
__device__ static inline unsigned int __clzll(unsigned long long int input) {
#ifdef NVCC_COMPAT
return input == 0 ? 64 : ( input == 0 ? -1 : __builtin_clzl(input) );
#else
return input == 0 ? -1 : __builtin_clzl(input);
#endif
}
__device__ static inline unsigned int __clz(int input) {
#ifdef NVCC_COMPAT
return input == 0 ? 32 : ( input > 0 ? __builtin_clz(input) : __builtin_clz(~input) );
#else
if (input == 0) return -1;
return input > 0 ? __builtin_clz(input) : __builtin_clz(~input);
#endif
}
__device__ static inline unsigned int __clzll(long long int input) {
#ifdef NVCC_COMPAT
return input == 0 ? 64 : input > 0 ? __builtin_clzl(input) : __builtin_clzl(~input);
#else
if (input == 0) return -1;
return input > 0 ? __builtin_clzl(input) : __builtin_clzl(~input);
#endif
__device__ static inline int __clzll(long long int input) {
return __ockl_clz_u64((ulong)input);
}
__device__ static inline unsigned int __ffs(unsigned int input) {
#ifdef NVCC_COMPAT
return ( input == 0 ? -1 : __builtin_ctz(input) ) + 1;
#else
return input == 0 ? -1 : __builtin_ctz(input);
#endif
}
__device__ static inline unsigned int __ffsll(unsigned long long int input) {
#ifdef NVCC_COMPAT
return ( input == 0 ? -1 : __builtin_ctzl(input) ) + 1;
#else
return input == 0 ? -1 : __builtin_ctzl(input);
#endif
}
__device__ static inline unsigned int __ffs(int input) {
#ifdef NVCC_COMPAT
return ( input == 0 ? -1 : __builtin_ctz(input) ) + 1;
#else
return input == 0 ? -1 : __builtin_ctz(input);
#endif
}
__device__ static inline unsigned int __ffsll(long long int input) {
#ifdef NVCC_COMPAT
return ( input == 0 ? -1 : __builtin_ctzl(input) ) + 1;
#else
return input == 0 ? -1 : __builtin_ctzl(input);
#endif
}
__device__ static inline unsigned int __brev(unsigned int input) { return __llvm_bitrev_b32(input); }
__device__ static inline unsigned int __brev(unsigned int input) {
return __llvm_bitrev_b32(input);
}
__device__ static inline unsigned long long int __brevll(unsigned long long int input) {
return __llvm_bitrev_b64(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) {
uint32_t offset = src1 & 31;
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;
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;
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 = (1 << width) - 1;
return ((src0 & ~(mask << offset)) | ((src1 & mask) << offset));
}
__device__ static unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int s);
__device__ static unsigned int __hadd(int x, int y);
__device__ static int __mul24(int x, int y);
@@ -606,7 +593,7 @@ __device__ static inline double __longlong_as_double(long long int x) {
double tmp;
__builtin_memcpy(&tmp, &x, sizeof(tmp));
return x;
return tmp;
}
__device__ static inline double __uint2double_rn(int x) { return (double)x; }
@@ -645,16 +632,43 @@ __device__ static inline float __ull2float_rz(unsigned long long int x) { return
#ifdef __HCC_OR_HIP_CLANG__
// Clock functions
__device__ long long int __clock64();
__device__ long long int __clock();
__device__ long long int clock64();
__device__ long long int clock();
// hip.amdgcn.bc - named sync
__device__ void __named_sync(int a, int b);
#ifdef __HIP_DEVICE_COMPILE__
// Clock functions
__device__
inline
long long int __clock64() { return (long long int) __builtin_amdgcn_s_memrealtime(); }
#if __HCC__
extern "C" uint64_t __clock_u64() __HC__;
#endif
__device__
inline
long long int __clock() { return (long long int) __builtin_amdgcn_s_memrealtime(); }
inline __attribute((always_inline))
long long int __clock64() {
// ToDo: Unify HCC and HIP implementation.
#if __HCC__
return (long long int) __clock_u64();
#else
return (long long int) __builtin_amdgcn_s_memrealtime();
#endif
}
__device__
inline __attribute((always_inline))
long long int __clock() { return __clock64(); }
__device__
inline __attribute__((always_inline))
long long int clock64() { return __clock64(); }
__device__
inline __attribute__((always_inline))
long long int clock() { return __clock(); }
// hip.amdgcn.bc - named sync
__device__
@@ -673,14 +687,7 @@ int __all(int predicate) {
__device__
inline
int __any(int predicate) {
#ifdef NVCC_COMPAT
if (__ockl_wfany_i32(predicate) != 0)
return 1;
else
return 0;
#else
return __ockl_wfany_i32(predicate);
#endif
}
// XXX from llvm/include/llvm/IR/InstrTypes.h
include/hip/hcc_detail/device_library_decls.h
+10
Wyświetl plik
@@ -30,6 +30,11 @@ THE SOFTWARE.
#include "hip/hcc_detail/host_defines.h"
typedef unsigned char uchar;
typedef unsigned short ushort;
typedef unsigned int uint;
typedef unsigned long ulong;
extern "C" __device__ __attribute__((const)) bool __ockl_wfany_i32(int);
extern "C" __device__ __attribute__((const)) bool __ockl_wfall_i32(int);
extern "C" __device__ uint __ockl_activelane_u32(void);
@@ -40,6 +45,11 @@ extern "C" __device__ __attribute__((const)) uint __ockl_mul_hi_u32(uint, uint);
extern "C" __device__ __attribute__((const)) int __ockl_mul_hi_i32(int, int);
extern "C" __device__ __attribute__((const)) uint __ockl_sad_u32(uint, uint, uint);
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)) ulong __ockl_clz_u64(ulong);
extern "C" __device__ __attribute__((const)) float __ocml_floor_f32(float);
extern "C" __device__ __attribute__((const)) float __ocml_rint_f32(float);
extern "C" __device__ __attribute__((const)) float __ocml_ceil_f32(float);
include/hip/hcc_detail/functional_grid_launch.hpp
+44 -20
Wyświetl plik
@@ -51,30 +51,53 @@ inline T round_up_to_next_multiple_nonnegative(T x, T y) {
return tmp - tmp % y;
}
inline std::vector<std::uint8_t> make_kernarg() { return {}; }
inline std::vector<std::uint8_t> make_kernarg(std::vector<std::uint8_t> kernarg) { return kernarg; }
template <typename T>
inline std::vector<std::uint8_t> make_kernarg(std::vector<uint8_t> kernarg, T x) {
kernarg.resize(round_up_to_next_multiple_nonnegative(kernarg.size(), alignof(T)) + sizeof(T));
new (kernarg.data() + kernarg.size() - sizeof(T)) T{std::move(x)};
template <
std::size_t n,
typename... Ts,
typename std::enable_if<n == sizeof...(Ts)>::type* = nullptr>
inline std::vector<std::uint8_t> make_kernarg(
std::vector<std::uint8_t> kernarg, const std::tuple<Ts...>&) {
return kernarg;
}
template <typename T, typename... Ts>
inline std::vector<std::uint8_t> make_kernarg(std::vector<std::uint8_t> kernarg, T x, Ts... xs) {
return make_kernarg(make_kernarg(std::move(kernarg), std::move(x)), std::move(xs)...);
template <
std::size_t n,
typename... Ts,
typename std::enable_if<n != sizeof...(Ts)>::type* = nullptr>
inline std::vector<std::uint8_t> make_kernarg(
std::vector<std::uint8_t> kernarg, const std::tuple<Ts...>& formals) {
using T = typename std::tuple_element<n, std::tuple<Ts...>>::type;
static_assert(
!std::is_reference<T>{},
"A __global__ function cannot have a reference as one of its "
"arguments.");
#if defined(HIP_STRICT)
static_assert(
std::is_trivially_copyable<T>{},
"Only TriviallyCopyable types can be arguments to a __global__ "
"function");
#endif
kernarg.resize(round_up_to_next_multiple_nonnegative(
kernarg.size(), alignof(T)) + sizeof(T));
new (kernarg.data() + kernarg.size() - sizeof(T)) T{std::get<n>(formals)};
return make_kernarg<n + 1>(std::move(kernarg), formals);
}
template <typename... Ts>
inline std::vector<std::uint8_t> make_kernarg(Ts... xs) {
std::vector<std::uint8_t> kernarg;
kernarg.reserve(sizeof(std::tuple<Ts...>));
template <typename... Formals, typename... Actuals>
inline std::vector<std::uint8_t> make_kernarg(
void (*)(Formals...), std::tuple<Actuals...> actuals) {
static_assert(sizeof...(Formals) == sizeof...(Actuals),
"The count of formal arguments must match the count of actuals.");
return make_kernarg(std::move(kernarg), std::move(xs)...);
std::tuple<Formals...> to_formals{std::move(actuals)};
std::vector<std::uint8_t> kernarg;
kernarg.reserve(sizeof(to_formals));
return make_kernarg<0>(std::move(kernarg), to_formals);
}
void hipLaunchKernelGGLImpl(std::uintptr_t function_address, const dim3& numBlocks,
@@ -85,7 +108,8 @@ void hipLaunchKernelGGLImpl(std::uintptr_t function_address, const dim3& numBloc
template <typename... Args, typename F = void (*)(Args...)>
inline void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
std::uint32_t sharedMemBytes, hipStream_t stream, Args... args) {
auto kernarg = hip_impl::make_kernarg(std::move(args)...);
auto kernarg = hip_impl::make_kernarg(
kernel, std::tuple<Args...>{std::move(args)...});
std::size_t kernarg_size = kernarg.size();
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, kernarg.data(), HIP_LAUNCH_PARAM_BUFFER_SIZE,
@@ -100,4 +124,4 @@ inline void hipLaunchKernel(F kernel, const dim3& numBlocks, const dim3& dimBloc
std::uint32_t groupMemBytes, hipStream_t stream, Args... args) {
hipLaunchKernelGGL(kernel, numBlocks, dimBlocks, groupMemBytes, stream, hipLaunchParm{},
std::move(args)...);
}
}
include/hip/hcc_detail/hip_complex.h
+3 -3
Wyświetl plik
@@ -121,7 +121,7 @@ THE SOFTWARE.
return ret; \
}
#define MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ComplexT, T) \
__device__ __host__ ComplexT(T val) : x(val), y(val) {} \
explicit __device__ __host__ ComplexT(T val) : x(val), y(val) {} \
__device__ __host__ ComplexT(T val1, T val2) : x(val1), y(val2) {}
#endif
@@ -131,7 +131,7 @@ struct hipFloatComplex {
public:
typedef float value_type;
__device__ __host__ hipFloatComplex() : x(0.0f), y(0.0f) {}
__device__ __host__ hipFloatComplex(float x) : x(x), y(0.0f) {}
explicit __device__ __host__ hipFloatComplex(float x) : x(x), y(0.0f) {}
__device__ __host__ hipFloatComplex(float x, float y) : x(x), y(y) {}
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, unsigned short)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, signed short)
@@ -151,7 +151,7 @@ struct hipDoubleComplex {
public:
typedef double value_type;
__device__ __host__ hipDoubleComplex() : x(0.0f), y(0.0f) {}
__device__ __host__ hipDoubleComplex(double x) : x(x), y(0.0f) {}
explicit __device__ __host__ hipDoubleComplex(double x) : x(x), y(0.0f) {}
__device__ __host__ hipDoubleComplex(double x, double y) : x(x), y(y) {}
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, unsigned short)
MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, signed short)
include/hip/hcc_detail/hip_fp16.h
+13 -13
Wyświetl plik
@@ -635,37 +635,37 @@ THE SOFTWARE.
// TODO: rounding behaviour is not correct.
// float -> half | half2
inline
__device__
__device__ __host__
__half __float2half(float x)
{
return __half_raw{static_cast<_Float16>(x)};
}
inline
__device__
__device__ __host__
__half __float2half_rn(float x)
{
return __half_raw{static_cast<_Float16>(x)};
}
inline
__device__
__device__ __host__
__half __float2half_rz(float x)
{
return __half_raw{static_cast<_Float16>(x)};
}
inline
__device__
__device__ __host__
__half __float2half_rd(float x)
{
return __half_raw{static_cast<_Float16>(x)};
}
inline
__device__
__device__ __host__
__half __float2half_ru(float x)
{
return __half_raw{static_cast<_Float16>(x)};
}
inline
__device__
__device__ __host__
__half2 __float2half2_rn(float x)
{
return __half2_raw{
@@ -673,14 +673,14 @@ THE SOFTWARE.
static_cast<_Float16>(x), static_cast<_Float16>(x)}};
}
inline
__device__
__device__ __host__
__half2 __floats2half2_rn(float x, float y)
{
return __half2_raw{_Float16_2{
static_cast<_Float16>(x), static_cast<_Float16>(y)}};
}
inline
__device__
__device__ __host__
__half2 __float22half2_rn(float2 x)
{
return __floats2half2_rn(x.x, x.y);
@@ -688,25 +688,25 @@ THE SOFTWARE.
// half | half2 -> float
inline
__device__
__device__ __host__
float __half2float(__half x)
{
return static_cast<__half_raw>(x).data;
}
inline
__device__
__device__ __host__
float __low2float(__half2 x)
{
return static_cast<__half2_raw>(x).data.x;
}
inline
__device__
__device__ __host__
float __high2float(__half2 x)
{
return static_cast<__half2_raw>(x).data.y;
}
inline
__device__
__device__ __host__
float2 __half22float2(__half2 x)
{
return make_float2(
@@ -1633,4 +1633,4 @@ THE SOFTWARE.
#endif // defined(__cplusplus)
#elif defined(__GNUC__)
#include "hip_fp16_gcc.h"
#endif // !defined(__clang__) && defined(__GNUC__)
#endif // !defined(__clang__) && defined(__GNUC__)
include/hip/hcc_detail/hip_runtime.h
+52 -27
Wyświetl plik
@@ -57,8 +57,6 @@ THE SOFTWARE.
#if __HCC_OR_HIP_CLANG__
// Define NVCC_COMPAT for CUDA compatibility
#define NVCC_COMPAT
#define CUDA_SUCCESS hipSuccess
#include <hip/hip_runtime_api.h>
@@ -110,9 +108,9 @@ extern int HIP_TRACE_API;
#include <hip/hcc_detail/host_defines.h>
#include <hip/hcc_detail/device_functions.h>
#include <hip/hcc_detail/surface_functions.h>
#include <hip/hcc_detail/texture_functions.h>
#if __HCC__
#include <hip/hcc_detail/math_functions.h>
#include <hip/hcc_detail/texture_functions.h>
#endif // __HCC__
// TODO-HCC remove old definitions ; ~1602 hcc supports __HCC_ACCELERATOR__ define.
@@ -201,16 +199,6 @@ __device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask
#endif //__HIP_ARCH_GFX803__ == 1
__device__ inline static int min(int arg1, int arg2) {
return (arg1 < arg2) ? arg1 : arg2;
}
__device__ inline static int max(int arg1, int arg2) {
return (arg1 > arg2) ? arg1 : arg2;
}
__host__ inline static int min(int arg1, int arg2) { return std::min(arg1, arg2); }
__host__ inline static int max(int arg1, int arg2) { return std::max(arg1, arg2); }
#endif // __HCC_OR_HIP_CLANG__
#if defined __HCC__
@@ -266,19 +254,16 @@ extern "C" __device__ void* __hip_free(void* ptr);
static inline __device__ void* malloc(size_t size) { return __hip_malloc(size); }
static inline __device__ void* free(void* ptr) { return __hip_free(ptr); }
#ifdef __HCC_ACCELERATOR__
#ifdef HC_FEATURE_PRINTF
#if defined(__HCC_ACCELERATOR__) && defined(HC_FEATURE_PRINTF)
template <typename... All>
static inline __device__ void printf(const char* format, All... all) {
hc::printf(format, all...);
}
#else
#elif defined(__HCC_ACCELERATOR__) || __HIP__
template <typename... All>
static inline __device__ void printf(const char* format, All... all) {}
#endif
#endif
#endif //__HCC_OR_HIP_CLANG__
#ifdef __HCC__
@@ -348,15 +333,18 @@ extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, gri
typedef int hipLaunchParm;
#define hipLaunchKernel(kernelName, numblocks, numthreads, memperblock, streamId, ...) \
do { \
kernelName<<<numblocks, numthreads, memperblock, streamId>>>(0, ##__VA_ARGS__); \
} while (0)
template <typename... Args, typename F = void (*)(Args...)>
inline void hipLaunchKernelGGL(F kernelName, const dim3& numblocks, const dim3& numthreads,
unsigned memperblock, hipStream_t streamId, Args... args) {
kernelName<<<numblocks, numthreads, memperblock, streamId>>>(args...);
}
#define hipLaunchKernelGGL(kernelName, numblocks, numthreads, memperblock, streamId, ...) \
do { \
kernelName<<<numblocks, numthreads, memperblock, streamId>>>(__VA_ARGS__); \
} while (0)
template <typename... Args, typename F = void (*)(hipLaunchParm, Args...)>
inline void hipLaunchKernel(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
std::uint32_t groupMemBytes, hipStream_t stream, Args... args) {
hipLaunchKernelGGL(kernel, numBlocks, dimBlocks, groupMemBytes, stream, hipLaunchParm{},
std::move(args)...);
}
#include <hip/hip_runtime_api.h>
@@ -440,6 +428,34 @@ extern const __device__ __attribute__((weak)) __hip_builtin_gridDim_t gridDim;
#include <hip/hcc_detail/math_functions.h>
#if __HIP_HCC_COMPAT_MODE__
// Define HCC work item functions in terms of HIP builtin variables.
#pragma push_macro("__DEFINE_HCC_FUNC")
#define __DEFINE_HCC_FUNC(hc_fun,hip_var) \
inline __device__ __attribute__((always_inline)) uint hc_get_##hc_fun(uint i) { \
if (i==0) \
return hip_var.x; \
else if(i==1) \
return hip_var.y; \
else \
return hip_var.z; \
}
__DEFINE_HCC_FUNC(workitem_id, threadIdx)
__DEFINE_HCC_FUNC(group_id, blockIdx)
__DEFINE_HCC_FUNC(group_size, blockDim)
__DEFINE_HCC_FUNC(num_groups, gridDim)
#pragma pop_macro("__DEFINE_HCC_FUNC")
extern "C" __device__ __attribute__((const)) size_t __ockl_get_global_id(uint);
inline __device__ __attribute__((always_inline)) uint
hc_get_workitem_absolute_id(int dim)
{
return (uint)__ockl_get_global_id(dim);
}
#endif
// Support std::complex.
#pragma push_macro("__CUDA__")
#define __CUDA__
@@ -447,11 +463,20 @@ extern const __device__ __attribute__((weak)) __hip_builtin_gridDim_t gridDim;
#include <__clang_cuda_complex_builtins.h>
#include <cuda_wrappers/algorithm>
#include <cuda_wrappers/complex>
#include <cuda_wrappers/new>
#undef __CUDA__
#pragma pop_macro("__CUDA__")
#endif
hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
uint32_t localWorkSizeX, uint32_t localWorkSizeY,
uint32_t localWorkSizeZ, size_t sharedMemBytes,
hipStream_t hStream, void** kernelParams, void** extra,
hipEvent_t startEvent = nullptr,
hipEvent_t stopEvent = nullptr);
#endif // defined(__clang__) && defined(__HIP__)
#include <hip/hcc_detail/hip_memory.h>
include/hip/hcc_detail/host_defines.h
+2 -2
Wyświetl plik
@@ -52,7 +52,7 @@ THE SOFTWARE.
#endif // GENERIC_GRID_LAUNCH
#define __noinline__ __attribute__((noinline))
#define __forceinline__ __attribute__((always_inline))
#define __forceinline__ inline __attribute__((always_inline))
/*
@@ -71,7 +71,7 @@ THE SOFTWARE.
#define __constant__ __attribute__((constant))
#define __noinline__ __attribute__((noinline))
#define __forceinline__ __attribute__((always_inline))
#define __forceinline__ inline __attribute__((always_inline))
#else
include/hip/hcc_detail/math_functions.h
+82
Wyświetl plik
@@ -31,6 +31,12 @@ THE SOFTWARE.
#include <limits>
#include <stdint.h>
// HCC's own math functions should be included first, otherwise there will
// be conflicts when hip/math_functions.h is included before hip/hip_runtime.h.
#ifdef __HCC__
#include "kalmar_math.h"
#endif
#pragma push_macro("__DEVICE__")
#pragma push_macro("__RETURN_TYPE")
@@ -1298,6 +1304,66 @@ float func(float x, int y) \
}
__DEF_FLOAT_FUN2I(scalbn)
#if __HCC__
template<class T>
__DEVICE__ inline static T min(T arg1, T arg2) {
return (arg1 < arg2) ? arg1 : arg2;
}
__DEVICE__ inline static uint32_t min(uint32_t arg1, int32_t arg2) {
return min(arg1, (uint32_t) arg2);
}
/*__DEVICE__ inline static uint32_t min(int32_t arg1, uint32_t arg2) {
return min((uint32_t) arg1, arg2);
}
__DEVICE__ inline static uint64_t min(uint64_t arg1, int64_t arg2) {
return min(arg1, (uint64_t) arg2);
}
__DEVICE__ inline static uint64_t min(int64_t arg1, uint64_t arg2) {
return min((uint64_t) arg1, arg2);
}
__DEVICE__ inline static unsigned long long min(unsigned long long arg1, long long arg2) {
return min(arg1, (unsigned long long) arg2);
}
__DEVICE__ inline static unsigned long long min(long long arg1, unsigned long long arg2) {
return min((unsigned long long) arg1, arg2);
}*/
template<class T>
__DEVICE__ inline static T max(T arg1, T arg2) {
return (arg1 > arg2) ? arg1 : arg2;
}
__DEVICE__ inline static uint32_t max(uint32_t arg1, int32_t arg2) {
return max(arg1, (uint32_t) arg2);
}
__DEVICE__ inline static uint32_t max(int32_t arg1, uint32_t arg2) {
return max((uint32_t) arg1, arg2);
}
/*__DEVICE__ inline static uint64_t max(uint64_t arg1, int64_t arg2) {
return max(arg1, (uint64_t) arg2);
}
__DEVICE__ inline static uint64_t max(int64_t arg1, uint64_t arg2) {
return max((uint64_t) arg1, arg2);
}
__DEVICE__ inline static unsigned long long max(unsigned long long arg1, long long arg2) {
return max(arg1, (unsigned long long) arg2);
}
__DEVICE__ inline static unsigned long long max(long long arg1, unsigned long long arg2) {
return max((unsigned long long) arg1, arg2);
}*/
#else
__DEVICE__ inline static int min(int arg1, int arg2) {
return (arg1 < arg2) ? arg1 : arg2;
}
__DEVICE__ inline static int max(int arg1, int arg2) {
return (arg1 > arg2) ? arg1 : arg2;
}
__DEVICE__
inline
float max(float x, float y) {
@@ -1325,6 +1391,17 @@ double min(double x, double y) {
__HIP_OVERLOAD2(double, max)
__HIP_OVERLOAD2(double, min)
#endif
__host__ inline static int min(int arg1, int arg2) {
return std::min(arg1, arg2);
}
__host__ inline static int max(int arg1, int arg2) {
return std::max(arg1, arg2);
}
#pragma pop_macro("__DEF_FLOAT_FUN")
#pragma pop_macro("__DEF_FLOAT_FUN2")
#pragma pop_macro("__DEF_FLOAT_FUN2I")
@@ -1332,3 +1409,8 @@ __HIP_OVERLOAD2(double, min)
#pragma pop_macro("__HIP_OVERLOAD2")
#pragma pop_macro("__DEVICE__")
#pragma pop_macro("__RETURN_TYPE")
// For backward compatibility.
// There are HIP applications e.g. TensorFlow, expecting __HIP_ARCH_* macros
// defined after including math_functions.h.
#include <hip/hcc_detail/hip_runtime.h>
include/hip/hcc_detail/program_state.hpp
+4 -5
Wyświetl plik
@@ -93,12 +93,11 @@ public:
}
};
const std::unordered_map<hsa_agent_t, std::vector<hsa_executable_t>>& executables(
bool rebuild = false);
const std::unordered_map<hsa_agent_t, std::vector<hsa_executable_t>>& executables();
const std::unordered_map<std::uintptr_t, std::vector<std::pair<hsa_agent_t, Kernel_descriptor>>>&
functions(bool rebuild = false);
const std::unordered_map<std::uintptr_t, std::string>& function_names(bool rebuild = false);
std::unordered_map<std::string, void*>& globals(bool rebuild = false);
functions();
const std::unordered_map<std::uintptr_t, std::string>& function_names();
std::unordered_map<std::string, void*>& globals();
hsa_executable_t load_executable(const std::string& file, hsa_executable_t executable,
hsa_agent_t agent);
include/hip/hcc_detail/texture_functions.h
+1344 -1309
Wyświetl plik
Plik diff jest za duży Load Diff
include/hip/math_functions.h
+1
Wyświetl plik
@@ -27,6 +27,7 @@ THE SOFTWARE.
// paths to provide a consistent include env and avoid "missing symbol" errors that only appears
// on NVCC path:
#include <hip/hip_common.h>
#if defined(__HIP_PLATFORM_HCC__) && !defined(__HIP_PLATFORM_NVCC__)
#include <hip/hcc_detail/math_functions.h>