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

Re-sync with upstream.

Browse Source
This commit is contained in:
Alex Voicu
2018-06-29 00:16:34 +01:00
parent 3c83e047df b326b8a0a6
commit fdb3e3e4a6
7 changed files with 82 additions and 48 deletions
Download Patch File Download Diff File Expand all files Collapse all files
include/hip/hcc_detail/hip_complex.h
+53 -5
View File
@@ -24,12 +24,34 @@ THE SOFTWARE.
#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H
#include "hip/hcc_detail/hip_vector_types.h"
#include <math.h>
// TODO: Clang has a bug which allows device functions to call std functions
// when std functions are introduced into default namespace by using statement.
// math.h may be included after this bug is fixed.
#if __cplusplus
#include <cmath>
#else
#include "math.h"
#endif
#if __cplusplus
#define MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(type, type1) \
__device__ __host__ type(type1 val) : x(val), y(val) {} \
__device__ __host__ type(type1 val1, type1 val2) : x(val1), y(val2) {}
#define COMPLEX_NEG_OP_OVERLOAD(type) \
__device__ __host__ static inline type operator-(const type& op) { \
type ret; \
ret.x = -op.x; \
ret.y = -op.y; \
return ret; \
}
#define COMPLEX_EQ_OP_OVERLOAD(type) \
__device__ __host__ static inline bool operator==(const type& lhs, const type& rhs) { \
return lhs.x == rhs.x && lhs.y == rhs.y; \
}
#define COMPLEX_NE_OP_OVERLOAD(type) \
__device__ __host__ static inline bool operator!=(const type& lhs, const type& rhs) { \
return !(lhs == rhs); \
}
#define COMPLEX_ADD_OP_OVERLOAD(type) \
__device__ __host__ static inline type operator+(const type& lhs, const type& rhs) { \
@@ -98,12 +120,16 @@ THE SOFTWARE.
ret.y = lhs.y * rhs; \
return ret; \
}
#define MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ComplexT, T) \
__device__ __host__ ComplexT(T val) : x(val), y(val) {} \
__device__ __host__ ComplexT(T val1, T val2) : x(val1), y(val2) {}
#endif
struct hipFloatComplex {
#ifdef __cplusplus
public:
typedef float value_type;
__device__ __host__ hipFloatComplex() : x(0.0f), y(0.0f) {}
__device__ __host__ hipFloatComplex(float x) : x(x), y(0.0f) {}
__device__ __host__ hipFloatComplex(float x, float y) : x(x), y(y) {}
@@ -123,6 +149,7 @@ struct hipFloatComplex {
struct hipDoubleComplex {
#ifdef __cplusplus
public:
typedef double value_type;
__device__ __host__ hipDoubleComplex() : x(0.0f), y(0.0f) {}
__device__ __host__ hipDoubleComplex(double x) : x(x), y(0.0f) {}
__device__ __host__ hipDoubleComplex(double x, double y) : x(x), y(y) {}
@@ -141,6 +168,9 @@ struct hipDoubleComplex {
#if __cplusplus
COMPLEX_NEG_OP_OVERLOAD(hipFloatComplex)
COMPLEX_EQ_OP_OVERLOAD(hipFloatComplex)
COMPLEX_NE_OP_OVERLOAD(hipFloatComplex)
COMPLEX_ADD_OP_OVERLOAD(hipFloatComplex)
COMPLEX_SUB_OP_OVERLOAD(hipFloatComplex)
COMPLEX_MUL_OP_OVERLOAD(hipFloatComplex)
@@ -160,6 +190,9 @@ COMPLEX_SCALAR_PRODUCT(hipFloatComplex, double)
COMPLEX_SCALAR_PRODUCT(hipFloatComplex, signed long long)
COMPLEX_SCALAR_PRODUCT(hipFloatComplex, unsigned long long)
COMPLEX_NEG_OP_OVERLOAD(hipDoubleComplex)
COMPLEX_EQ_OP_OVERLOAD(hipDoubleComplex)
COMPLEX_NE_OP_OVERLOAD(hipDoubleComplex)
COMPLEX_ADD_OP_OVERLOAD(hipDoubleComplex)
COMPLEX_SUB_OP_OVERLOAD(hipDoubleComplex)
COMPLEX_MUL_OP_OVERLOAD(hipDoubleComplex)
@@ -225,7 +258,6 @@ __device__ __host__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hi
__device__ __host__ static inline float hipCabsf(hipFloatComplex z) { return sqrtf(hipCsqabsf(z)); }
__device__ __host__ static inline double hipCreal(hipDoubleComplex z) { return z.x; }
__device__ __host__ static inline double hipCimag(hipDoubleComplex z) { return z.y; }
@@ -305,4 +337,20 @@ __device__ __host__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, h
return make_hipDoubleComplex(real, imag);
}
// Complex functions returning real numbers.
#define __DEFINE_HIP_COMPLEX_REAL_FUN(func, hipFun) \
__device__ __host__ inline float func(const hipFloatComplex& z) { return hipFun##f(z); } \
__device__ __host__ inline double func(const hipDoubleComplex& z) { return hipFun(z); }
__DEFINE_HIP_COMPLEX_REAL_FUN(abs, hipCabs)
__DEFINE_HIP_COMPLEX_REAL_FUN(real, hipCreal)
__DEFINE_HIP_COMPLEX_REAL_FUN(imag, hipCimag)
// Complex functions returning complex numbers.
#define __DEFINE_HIP_COMPLEX_FUN(func, hipFun) \
__device__ __host__ inline hipFloatComplex func(const hipFloatComplex& z) { return hipFun##f(z); } \
__device__ __host__ inline hipDoubleComplex func(const hipDoubleComplex& z) { return hipFun(z); }
__DEFINE_HIP_COMPLEX_FUN(conj, hipConj)
#endif
include/hip/hcc_detail/hip_fp16.h
+1 -1
View File
@@ -21,7 +21,7 @@ THE SOFTWARE.
*/
#pragma once
#include "hip/hcc_detail/host_defines.h"
#include <assert.h>
#if defined(__cplusplus)
#include <algorithm>
include/hip/hcc_detail/hip_runtime.h
+5 -5
View File
@@ -98,6 +98,9 @@ struct Empty_launch_parm {};
#include "grid_launch_GGL.hpp"
#endif // GENERIC_GRID_LAUNCH
#endif // HCC
#if __HCC_OR_HIP_CLANG__
extern int HIP_TRACE_API;
#ifdef __cplusplus
@@ -107,15 +110,16 @@ extern int HIP_TRACE_API;
#include <hip/hcc_detail/host_defines.h>
#include <hip/hcc_detail/math_functions.h>
#include <hip/hcc_detail/device_functions.h>
#if __HCC__
#include <hip/hcc_detail/texture_functions.h>
#include <hip/hcc_detail/surface_functions.h>
#endif // __HCC__
// TODO-HCC remove old definitions ; ~1602 hcc supports __HCC_ACCELERATOR__ define.
#if defined(__KALMAR_ACCELERATOR__) && !defined(__HCC_ACCELERATOR__)
#define __HCC_ACCELERATOR__ __KALMAR_ACCELERATOR__
#endif
// TODO-HCC add a dummy implementation of assert, need to replace with a proper kernel exit call.
#if __HIP_DEVICE_COMPILE__ == 1
#undef assert
@@ -180,10 +184,6 @@ extern int HIP_TRACE_API;
#define __HCC_C__
#endif
#endif // defined __HCC__
#if __HCC_OR_HIP_CLANG__
// TODO - hipify-clang - change to use the function call.
//#define warpSize hc::__wavesize()
static constexpr int warpSize = 64;
include/hip/hcc_detail/math_functions.h
+9 -9
View File
@@ -508,19 +508,19 @@ inline
float __fadd_rz(float x, float y) { return __ocml_add_rtz_f32(x, y); }
__device__
inline
float __fdiv_rd(float x, float y) { return __ocml_div_rtp_f32(x, y); }
float __fdiv_rd(float x, float y) { return x / y; }
__device__
inline
float __fdiv_rn(float x, float y) { return __ocml_div_rte_f32(x, y); }
float __fdiv_rn(float x, float y) { return x / y; }
__device__
inline
float __fdiv_ru(float x, float y) { return __ocml_div_rtn_f32(x, y); }
float __fdiv_ru(float x, float y) { return x / y; }
__device__
inline
float __fdiv_rz(float x, float y) { return __ocml_div_rtz_f32(x, y); }
float __fdiv_rz(float x, float y) { return x / y; }
__device__
inline
float __fdividef(float x, float y) { return __ocml_div_rte_f32(x, y); }
float __fdividef(float x, float y) { return x / y; }
__device__
inline
float __fmaf_rd(float x, float y, float z)
@@ -1028,16 +1028,16 @@ inline
double __dadd_rz(double x, double y) { return __ocml_add_rtz_f64(x, y); }
__device__
inline
double __ddiv_rd(double x, double y) { return __ocml_div_rtp_f64(x, y); }
double __ddiv_rd(double x, double y) { return x / y; }
__device__
inline
double __ddiv_rn(double x, double y) { return __ocml_div_rte_f64(x, y); }
double __ddiv_rn(double x, double y) { return x / y; }
__device__
inline
double __ddiv_ru(double x, double y) { return __ocml_div_rtn_f64(x, y); }
double __ddiv_ru(double x, double y) { return x / y; }
__device__
inline
double __ddiv_rz(double x, double y) { return __ocml_div_rtz_f64(x, y); }
double __ddiv_rz(double x, double y) { return x / y; }
__device__
inline
double __dmul_rd(double x, double y) { return __ocml_mul_rtp_f64(x, y); }
include/hip/hcc_detail/math_fwd.h
-26
View File
@@ -69,8 +69,6 @@ float __ocml_cosh_f32(float);
__device__
float __ocml_cospi_f32(float);
__device__
float __ocml_div_rtz_f32(float, float);
__device__
float __ocml_i0_f32(float);
__device__
float __ocml_i1_f32(float);
@@ -290,18 +288,6 @@ __attribute__((const))
float __ocml_mul_rtz_f32(float, float);
__device__
__attribute__((const))
float __ocml_div_rte_f32(float, float);
__device__
__attribute__((const))
float __ocml_div_rtn_f32(float, float);
__device__
__attribute__((const))
float __ocml_div_rtp_f32(float, float);
__device__
__attribute__((const))
float __ocml_div_rtz_f32(float, float);
__device__
__attribute__((const))
float __ocml_sqrt_rte_f32(float);
__device__
__attribute__((const))
@@ -598,18 +584,6 @@ __attribute__((const))
double __ocml_mul_rtz_f64(double, double);
__device__
__attribute__((const))
double __ocml_div_rte_f64(double, double);
__device__
__attribute__((const))
double __ocml_div_rtn_f64(double, double);
__device__
__attribute__((const))
double __ocml_div_rtp_f64(double, double);
__device__
__attribute__((const))
double __ocml_div_rtz_f64(double, double);
__device__
__attribute__((const))
double __ocml_sqrt_rte_f64(double);
__device__
__attribute__((const))