AI/rocm-systems
2
0
انشعاب 0
کد مسائل تقاضاهای واکشی Actions Packages پروژه‌ها انتشارها دانشنامه فعالیت

Merge branch 'master' of https://github.com/ROCm-Developer-Tools/HIP into feature_get_alignment_and_size_from_metadata

فهرست منبع
This commit is contained in:
Alex Voicu
2018-11-08 06:53:06 -08:00
والد 1c525d7222 e510be9395
کامیت 392e2ba1ba
12فایلهای تغییر یافته به همراه236 افزوده شده و 113 حذف شده
دانلود پرونده وصله دانلود فایل تغییرات diff Expand all files Collapse all files
hipamd/include/hip/hcc_detail/hip_memory.h
+4 -5
مشاهده پرونده
@@ -27,6 +27,8 @@ THE SOFTWARE.
// HIP heap is implemented as a global array with fixed size. Users may define
// __HIP_SIZE_OF_PAGE and __HIP_NUM_PAGES to have a larger heap.
#if __HCC__ || __HIP__
// Size of page in bytes.
#ifndef __HIP_SIZE_OF_PAGE
#define __HIP_SIZE_OF_PAGE 64
@@ -39,11 +41,8 @@ THE SOFTWARE.
#define __HIP_SIZE_OF_HEAP (__HIP_NUM_PAGES * __HIP_SIZE_OF_PAGE)
#if __HCC__ || __HIP__
__attribute__((weak)) __device__ char __hip_device_heap[__HIP_SIZE_OF_HEAP];
__attribute__((weak)) __device__
uint32_t __hip_device_page_flag[__HIP_NUM_PAGES];
extern __device__ char __hip_device_heap[];
extern __device__ uint32_t __hip_device_page_flag[];
extern "C" inline __device__ void* __hip_malloc(size_t size) {
char* heap = (char*)__hip_device_heap;
hipamd/include/hip/hcc_detail/math_functions.h
+38 -32
مشاهده پرونده
@@ -514,38 +514,41 @@ float __exp10f(float x) { return __ocml_exp10_f32(x); }
__DEVICE__
inline
float __expf(float x) { return __ocml_exp_f32(x); }
#if defined OCML_BASIC_ROUNDED_OPERATIONS
__DEVICE__
inline
float __fadd_rd(float x, float y) { return __ocml_add_rtp_f32(x, y); }
float __fadd_rd(float x, float y) { return __ocml_add_rtn_f32(x, y); }
__DEVICE__
inline
float __fadd_rn(float x, float y) { return __ocml_add_rte_f32(x, y); }
__DEVICE__
inline
float __fadd_ru(float x, float y) { return __ocml_add_rtn_f32(x, y); }
float __fadd_ru(float x, float y) { return __ocml_add_rtp_f32(x, y); }
__DEVICE__
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 x / y; }
float __fdiv_rd(float x, float y) { return __ocml_div_rtn_f32(x, y); }
__DEVICE__
inline
float __fdiv_rn(float x, float y) { return x / y; }
float __fdiv_rn(float x, float y) { return __ocml_div_rte_f32(x, y); }
__DEVICE__
inline
float __fdiv_ru(float x, float y) { return x / y; }
float __fdiv_ru(float x, float y) { return __ocml_div_rtp_f32(x, y); }
__DEVICE__
inline
float __fdiv_rz(float x, float y) { return x / y; }
float __fdiv_rz(float x, float y) { return __ocml_div_rtz_f32(x, y); }
#endif
__DEVICE__
inline
float __fdividef(float x, float y) { return x / y; }
#if defined OCML_BASIC_ROUNDED_OPERATIONS
__DEVICE__
inline
float __fmaf_rd(float x, float y, float z)
{
return __ocml_fma_rtp_f32(x, y, z);
return __ocml_fma_rtn_f32(x, y, z);
}
__DEVICE__
inline
@@ -557,7 +560,7 @@ __DEVICE__
inline
float __fmaf_ru(float x, float y, float z)
{
return __ocml_fma_rtn_f32(x, y, z);
return __ocml_fma_rtp_f32(x, y, z);
}
__DEVICE__
inline
@@ -567,13 +570,13 @@ float __fmaf_rz(float x, float y, float z)
}
__DEVICE__
inline
float __fmul_rd(float x, float y) { return __ocml_mul_rtp_f32(x, y); }
float __fmul_rd(float x, float y) { return __ocml_mul_rtn_f32(x, y); }
__DEVICE__
inline
float __fmul_rn(float x, float y) { return __ocml_mul_rte_f32(x, y); }
__DEVICE__
inline
float __fmul_ru(float x, float y) { return __ocml_mul_rtn_f32(x, y); }
float __fmul_ru(float x, float y) { return __ocml_mul_rtp_f32(x, y); }
__DEVICE__
inline
float __fmul_rz(float x, float y) { return __ocml_mul_rtz_f32(x, y); }
@@ -594,28 +597,29 @@ inline
float __frsqrt_rn(float x) { return __llvm_amdgcn_rsq_f32(x); }
__DEVICE__
inline
float __fsqrt_rd(float x) { return __ocml_sqrt_f32(x); }
float __fsqrt_rd(float x) { return __ocml_sqrt_rtn_f32(x); }
__DEVICE__
inline
float __fsqrt_rn(float x) { return __ocml_sqrt_f32(x); }
float __fsqrt_rn(float x) { return __ocml_sqrt_rte_f32(x); }
__DEVICE__
inline
float __fsqrt_ru(float x) { return __ocml_sqrt_f32(x); }
float __fsqrt_ru(float x) { return __ocml_sqrt_rtp_f32(x); }
__DEVICE__
inline
float __fsqrt_rz(float x) { return __ocml_sqrt_f32(x); }
float __fsqrt_rz(float x) { return __ocml_sqrt_rtz_f32(x); }
__DEVICE__
inline
float __fsub_rd(float x, float y) { return __ocml_sub_rtp_f32(x, y); }
float __fsub_rd(float x, float y) { return __ocml_sub_rtn_f32(x, y); }
__DEVICE__
inline
float __fsub_rn(float x, float y) { return __ocml_sub_rte_f32(x, y); }
__DEVICE__
inline
float __fsub_ru(float x, float y) { return __ocml_sub_rtn_f32(x, y); }
float __fsub_ru(float x, float y) { return __ocml_sub_rtp_f32(x, y); }
__DEVICE__
inline
float __fsub_rz(float x, float y) { return __ocml_sub_rtz_f32(x, y); }
#endif
__DEVICE__
inline
float __log10f(float x) { return __ocml_log10_f32(x); }
@@ -1034,39 +1038,40 @@ double yn(int n, double x)
}
// BEGIN INTRINSICS
#if defined OCML_BASIC_ROUNDED_OPERATIONS
__DEVICE__
inline
double __dadd_rd(double x, double y) { return __ocml_add_rtp_f64(x, y); }
double __dadd_rd(double x, double y) { return __ocml_add_rtn_f64(x, y); }
__DEVICE__
inline
double __dadd_rn(double x, double y) { return __ocml_add_rte_f64(x, y); }
__DEVICE__
inline
double __dadd_ru(double x, double y) { return __ocml_add_rtn_f64(x, y); }
double __dadd_ru(double x, double y) { return __ocml_add_rtp_f64(x, y); }
__DEVICE__
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 x / y; }
double __ddiv_rd(double x, double y) { return __ocml_div_rtn_f64(x, y); }
__DEVICE__
inline
double __ddiv_rn(double x, double y) { return x / y; }
double __ddiv_rn(double x, double y) { return __ocml_div_rte_f64(x, y); }
__DEVICE__
inline
double __ddiv_ru(double x, double y) { return x / y; }
double __ddiv_ru(double x, double y) { return __ocml_div_rtp_f64(x, y); }
__DEVICE__
inline
double __ddiv_rz(double x, double y) { return x / y; }
double __ddiv_rz(double x, double y) { return __ocml_div_rtz_f64(x, y); }
__DEVICE__
inline
double __dmul_rd(double x, double y) { return __ocml_mul_rtp_f64(x, y); }
double __dmul_rd(double x, double y) { return __ocml_mul_rtn_f64(x, y); }
__DEVICE__
inline
double __dmul_rn(double x, double y) { return __ocml_mul_rte_f64(x, y); }
__DEVICE__
inline
double __dmul_ru(double x, double y) { return __ocml_mul_rtn_f64(x, y); }
double __dmul_ru(double x, double y) { return __ocml_mul_rtp_f64(x, y); }
__DEVICE__
inline
double __dmul_rz(double x, double y) { return __ocml_mul_rtz_f64(x, y); }
@@ -1084,25 +1089,25 @@ inline
double __drcp_rz(double x) { return __llvm_amdgcn_rcp_f64(x); }
__DEVICE__
inline
double __dsqrt_rd(double x) { return __ocml_sqrt_f64(x); }
double __dsqrt_rd(double x) { return __ocml_sqrt_rtn_f64(x); }
__DEVICE__
inline
double __dsqrt_rn(double x) { return __ocml_sqrt_f64(x); }
double __dsqrt_rn(double x) { return __ocml_sqrt_rte_f64(x); }
__DEVICE__
inline
double __dsqrt_ru(double x) { return __ocml_sqrt_f64(x); }
double __dsqrt_ru(double x) { return __ocml_sqrt_rtp_f64(x); }
__DEVICE__
inline
double __dsqrt_rz(double x) { return __ocml_sqrt_f64(x); }
double __dsqrt_rz(double x) { return __ocml_sqrt_rtz_f64(x); }
__DEVICE__
inline
double __dsub_rd(double x, double y) { return __ocml_sub_rtp_f64(x, y); }
double __dsub_rd(double x, double y) { return __ocml_sub_rtn_f64(x, y); }
__DEVICE__
inline
double __dsub_rn(double x, double y) { return __ocml_sub_rte_f64(x, y); }
__DEVICE__
inline
double __dsub_ru(double x, double y) { return __ocml_sub_rtn_f64(x, y); }
double __dsub_ru(double x, double y) { return __ocml_sub_rtp_f64(x, y); }
__DEVICE__
inline
double __dsub_rz(double x, double y) { return __ocml_sub_rtz_f64(x, y); }
@@ -1110,7 +1115,7 @@ __DEVICE__
inline
double __fma_rd(double x, double y, double z)
{
return __ocml_fma_rtp_f64(x, y, z);
return __ocml_fma_rtn_f64(x, y, z);
}
__DEVICE__
inline
@@ -1122,7 +1127,7 @@ __DEVICE__
inline
double __fma_ru(double x, double y, double z)
{
return __ocml_fma_rtn_f64(x, y, z);
return __ocml_fma_rtp_f64(x, y, z);
}
__DEVICE__
inline
@@ -1130,6 +1135,7 @@ double __fma_rz(double x, double y, double z)
{
return __ocml_fma_rtz_f64(x, y, z);
}
#endif
// END INTRINSICS
// END DOUBLE
hipamd/include/hip/hcc_detail/math_fwd.h
+49 -1
مشاهده پرونده
@@ -288,6 +288,30 @@ __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, float);
__device__
__attribute__((const))
float __ocml_sqrt_rtn_f32(float, float);
__device__
__attribute__((const))
float __ocml_sqrt_rtp_f32(float, float);
__device__
__attribute__((const))
float __ocml_sqrt_rtz_f32(float, float);
__device__
__attribute__((const))
float __ocml_fma_rte_f32(float, float, float);
__device__
__attribute__((const))
@@ -572,6 +596,30 @@ __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, double);
__device__
__attribute__((const))
double __ocml_sqrt_rtn_f64(double, double);
__device__
__attribute__((const))
double __ocml_sqrt_rtp_f64(double, double);
__device__
__attribute__((const))
double __ocml_sqrt_rtz_f64(double, double);
__device__
__attribute__((const))
double __ocml_fma_rte_f64(double, double, double);
__device__
__attribute__((const))
@@ -594,4 +642,4 @@ double __llvm_amdgcn_rsq_f64(double) __asm("llvm.amdgcn.rsq.f64");
#if defined(__cplusplus)
} // extern "C"
#endif
#endif