diff --git a/hipamd/include/hip/hcc_detail/grid_launch.h b/hipamd/include/hip/hcc_detail/grid_launch.h new file mode 100644 index 0000000000..f91d23341a --- /dev/null +++ b/hipamd/include/hip/hcc_detail/grid_launch.h @@ -0,0 +1,69 @@ +#pragma once + +#include + +#include + +#define GRID_LAUNCH_VERSION 20 + +// Extern definitions +namespace hc{ +class completion_future; +class accelerator_view; +} + + +// 3 dim structure for groups and grids. +typedef struct gl_dim3 +{ + int x,y,z; + gl_dim3(uint32_t _x=1, uint32_t _y=1, uint32_t _z=1) : x(_x), y(_y), z(_z) {}; +} gl_dim3; + +typedef enum gl_barrier_bit { + barrier_bit_queue_default, + barrier_bit_none, + barrier_bit_wait, +} gl_barrier_bit; + + +// grid_launch_parm contains information used to launch the kernel. +typedef struct grid_launch_parm +{ + //! Grid dimensions + gl_dim3 grid_dim; + + //! Group dimensions + gl_dim3 group_dim;; + + //! Amount of dynamic group memory to use with the kernel launch. + //! This memory is in addition to the amount used statically in the kernel. + unsigned int dynamic_group_mem_bytes;; + + //! Control setting of barrier bit on per-packet basis: + //! See gl_barrier_bit description. + //! Placeholder, is not used to control packet dispatch yet + enum gl_barrier_bit barrier_bit; + + //! Value of packet fences to apply to launch. + //! The correspond to the value of bits 9:14 in the AQL packet, + //! see HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE and hsa_fence_scope_t. + //! Set to -1 for conservative defaults. + //! Placeholder, is not used to control packet dispatch yet + unsigned int launch_fence; + + //! Pointer to the accelerator_view where the kernel should execute. + //! If NULL, the default view on the default accelerator is used. + hc::accelerator_view *av; + + //! Pointe to the completion_future used to track the status of the command. + //! If NULL, the command does not write status. In this case, + //! synchronization can be enforced with queue-level waits or + //! waiting on younger commands. + hc::completion_future *cf; + + grid_launch_parm() = default; +} grid_launch_parm; + + +extern void init_grid_launch(grid_launch_parm *gl); diff --git a/hipamd/include/hip/hcc_detail/grid_launch.hpp b/hipamd/include/hip/hcc_detail/grid_launch.hpp new file mode 100644 index 0000000000..04ce7e0366 --- /dev/null +++ b/hipamd/include/hip/hcc_detail/grid_launch.hpp @@ -0,0 +1,50 @@ +#pragma once + +#include "grid_launch.h" +#include "hc.hpp" + +class grid_launch_parm_cxx : public grid_launch_parm +{ +public: + grid_launch_parm_cxx() = default; + + // customized serialization: don't need av and cf in kernel + __attribute__((annotate("serialize"))) + void __cxxamp_serialize(Kalmar::Serialize& s) const { + s.Append(sizeof(int), &grid_dim.x); + s.Append(sizeof(int), &grid_dim.y); + s.Append(sizeof(int), &grid_dim.z); + s.Append(sizeof(int), &group_dim.x); + s.Append(sizeof(int), &group_dim.y); + s.Append(sizeof(int), &group_dim.z); + } + + __attribute__((annotate("user_deserialize"))) + grid_launch_parm_cxx(int grid_dim_x, int grid_dim_y, int grid_dim_z, + int group_dim_x, int group_dim_y, int group_dim_z) { + grid_dim.x = grid_dim_x; + grid_dim.y = grid_dim_y; + grid_dim.z = grid_dim_z; + group_dim.x = group_dim_x; + group_dim.y = group_dim_y; + group_dim.z = group_dim_z; + } +}; + + +extern inline void grid_launch_init(grid_launch_parm *lp) { + lp->grid_dim.x = lp->grid_dim.y = lp->grid_dim.z = 1; + + lp->group_dim.x = lp->group_dim.y = lp->group_dim.z = 1; + + lp->dynamic_group_mem_bytes = 0; + + lp->barrier_bit = barrier_bit_queue_default; + lp->launch_fence = -1; + + // TODO - set to NULL? + static hc::accelerator_view av = hc::accelerator().get_default_view(); + lp->av = &av; + lp->cf = NULL; +} + diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h index 7a6aba5ab3..9759a9bdf8 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -71,7 +71,7 @@ THE SOFTWARE. //--- // Remainder of this file only compiles with HCC #if defined __HCC__ -#include +#include "grid_launch.h" #include "hc_printf.hpp" // TODO-HCC-GL - change this to typedef. // typedef grid_launch_parm hipLaunchParm ; @@ -109,9 +109,6 @@ extern int HIP_TRACE_API; #include #include #include -#if __HCC__ -#include -#endif // __HCC__ // TODO-HCC remove old definitions ; ~1602 hcc supports __HCC_ACCELERATOR__ define. #if defined(__KALMAR_ACCELERATOR__) && !defined(__HCC_ACCELERATOR__) diff --git a/hipamd/include/hip/hcc_detail/math_functions.h b/hipamd/include/hip/hcc_detail/math_functions.h index 6c8510fcbb..c7a4d0e88f 100644 --- a/hipamd/include/hip/hcc_detail/math_functions.h +++ b/hipamd/include/hip/hcc_detail/math_functions.h @@ -26,17 +26,12 @@ THE SOFTWARE. #include +#include #include #include #include #include -// 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") @@ -592,16 +587,16 @@ inline float __frsqrt_rn(float x) { return __llvm_amdgcn_rsq_f32(x); } __DEVICE__ inline -float __fsqrt_rd(float x) { return __ocml_sqrt_rtp_f32(x); } +float __fsqrt_rd(float x) { return __ocml_sqrt_f32(x); } __DEVICE__ inline -float __fsqrt_rn(float x) { return __ocml_sqrt_rte_f32(x); } +float __fsqrt_rn(float x) { return __ocml_sqrt_f32(x); } __DEVICE__ inline -float __fsqrt_ru(float x) { return __ocml_sqrt_rtn_f32(x); } +float __fsqrt_ru(float x) { return __ocml_sqrt_f32(x); } __DEVICE__ inline -float __fsqrt_rz(float x) { return __ocml_sqrt_rtz_f32(x); } +float __fsqrt_rz(float x) { return __ocml_sqrt_f32(x); } __DEVICE__ inline float __fsub_rd(float x, float y) { return __ocml_sub_rtp_f32(x, y); } @@ -1082,16 +1077,16 @@ inline double __drcp_rz(double x) { return __llvm_amdgcn_rcp_f64(x); } __DEVICE__ inline -double __dsqrt_rd(double x) { return __ocml_sqrt_rtp_f64(x); } +double __dsqrt_rd(double x) { return __ocml_sqrt_f64(x); } __DEVICE__ inline -double __dsqrt_rn(double x) { return __ocml_sqrt_rte_f64(x); } +double __dsqrt_rn(double x) { return __ocml_sqrt_f64(x); } __DEVICE__ inline -double __dsqrt_ru(double x) { return __ocml_sqrt_rtn_f64(x); } +double __dsqrt_ru(double x) { return __ocml_sqrt_f64(x); } __DEVICE__ inline -double __dsqrt_rz(double x) { return __ocml_sqrt_rtz_f64(x); } +double __dsqrt_rz(double x) { return __ocml_sqrt_f64(x); } __DEVICE__ inline double __dsub_rd(double x, double y) { return __ocml_sub_rtp_f64(x, y); } diff --git a/hipamd/include/hip/hcc_detail/math_fwd.h b/hipamd/include/hip/hcc_detail/math_fwd.h index 102714e117..404c2f81d5 100644 --- a/hipamd/include/hip/hcc_detail/math_fwd.h +++ b/hipamd/include/hip/hcc_detail/math_fwd.h @@ -288,18 +288,6 @@ __attribute__((const)) float __ocml_mul_rtz_f32(float, float); __device__ __attribute__((const)) -float __ocml_sqrt_rte_f32(float); -__device__ -__attribute__((const)) -float __ocml_sqrt_rtn_f32(float); -__device__ -__attribute__((const)) -float __ocml_sqrt_rtp_f32(float); -__device__ -__attribute__((const)) -float __ocml_sqrt_rtz_f32(float); -__device__ -__attribute__((const)) float __ocml_fma_rte_f32(float, float, float); __device__ __attribute__((const)) @@ -584,18 +572,6 @@ __attribute__((const)) double __ocml_mul_rtz_f64(double, double); __device__ __attribute__((const)) -double __ocml_sqrt_rte_f64(double); -__device__ -__attribute__((const)) -double __ocml_sqrt_rtn_f64(double); -__device__ -__attribute__((const)) -double __ocml_sqrt_rtp_f64(double); -__device__ -__attribute__((const)) -double __ocml_sqrt_rtz_f64(double); -__device__ -__attribute__((const)) double __ocml_fma_rte_f64(double, double, double); __device__ __attribute__((const)) diff --git a/hipamd/src/device_util.cpp b/hipamd/src/device_util.cpp index 87fbe0fcbc..2e78dea009 100644 --- a/hipamd/src/device_util.cpp +++ b/hipamd/src/device_util.cpp @@ -21,10 +21,9 @@ THE SOFTWARE. */ #include -#include -#include #include "device_util.h" #include "hip/hcc_detail/device_functions.h" +#include "hip/hcc_detail/grid_launch.h" #include "hip/hip_runtime.h" #include diff --git a/hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp b/hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp index 76b7a5de7a..413f67d012 100644 --- a/hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp +++ b/hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp @@ -28,6 +28,7 @@ THE SOFTWARE. #include #include +#include #include "test_common.h" #pragma GCC diagnostic ignored "-Wall" diff --git a/hipamd/tests/src/deviceLib/hipFloatMath.cpp b/hipamd/tests/src/deviceLib/hipFloatMath.cpp index 157ea046ec..403a77f703 100644 --- a/hipamd/tests/src/deviceLib/hipFloatMath.cpp +++ b/hipamd/tests/src/deviceLib/hipFloatMath.cpp @@ -27,6 +27,7 @@ THE SOFTWARE. */ #include "test_common.h" +#include #include #define LEN 512 diff --git a/hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp b/hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp index 1e6a645b1a..d62a591257 100644 --- a/hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp +++ b/hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp @@ -31,6 +31,10 @@ THE SOFTWARE. #include #include "test_common.h" +#include + +using namespace std; + #pragma GCC diagnostic ignored "-Wall" #pragma clang diagnostic ignored "-Wunused-variable" diff --git a/hipamd/tests/src/deviceLib/hipMathFunctions.cpp b/hipamd/tests/src/deviceLib/hipMathFunctions.cpp index f89bdae149..3acafddcc0 100644 --- a/hipamd/tests/src/deviceLib/hipMathFunctions.cpp +++ b/hipamd/tests/src/deviceLib/hipMathFunctions.cpp @@ -27,6 +27,7 @@ THE SOFTWARE. */ #include "hip/hip_runtime.h" +#include "hip/math_functions.h" #include "test_common.h" #if __HIP_ARCH_GFX803__ || __HIP_ARCH_GFX900__ || __HIP_ARCH_GFX906__ diff --git a/hipamd/tests/src/deviceLib/hipSimpleAtomicsTest.cpp b/hipamd/tests/src/deviceLib/hipSimpleAtomicsTest.cpp index 129d7c1aa8..4890ed0f97 100644 --- a/hipamd/tests/src/deviceLib/hipSimpleAtomicsTest.cpp +++ b/hipamd/tests/src/deviceLib/hipSimpleAtomicsTest.cpp @@ -215,6 +215,7 @@ template< typename T, typename enable_if< is_same{} || is_same{}>::type* = nullptr> +__device__ void testKernelSub(T* g_odata) { // Atomic subtraction (final should be 0) atomicSub(&g_odata[1], 10); diff --git a/hipamd/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp b/hipamd/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp index dab153ed5e..c0d21aade8 100644 --- a/hipamd/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp +++ b/hipamd/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp @@ -27,6 +27,8 @@ THE SOFTWARE. #include #include +#include + #include "test_common.h" #pragma GCC diagnostic ignored "-Wall" diff --git a/hipamd/tests/src/deviceLib/hip_trig.cpp b/hipamd/tests/src/deviceLib/hip_trig.cpp index 3f200fcf8f..e9bc1dcf13 100644 --- a/hipamd/tests/src/deviceLib/hip_trig.cpp +++ b/hipamd/tests/src/deviceLib/hip_trig.cpp @@ -28,9 +28,11 @@ THE SOFTWARE. #include #include -#include #include "test_common.h" #include +#include + +#include #define LEN 512 #define SIZE LEN << 2 diff --git a/hipamd/tests/src/kernel/hipLanguageExtensions.cpp b/hipamd/tests/src/kernel/hipLanguageExtensions.cpp index 36b79965ca..2677d4da8e 100644 --- a/hipamd/tests/src/kernel/hipLanguageExtensions.cpp +++ b/hipamd/tests/src/kernel/hipLanguageExtensions.cpp @@ -31,7 +31,7 @@ THE SOFTWARE. #include #ifdef __HCC__ -#include + #include #endif // cudaA @@ -93,11 +93,7 @@ __global__ void vectorADD(const hipLaunchParm lp, T __restrict__* A_d, T* B_d, T int b = threadIdx.x; int c; - // TODO - move to HIP atomics when ready. - concurrency ::atomic_fetch_add(&c, b); - // Concurrency::atomic_add_unsigned (&x, a); - - // concurrency ::atomic_add_ (x, a); + atomicAdd(&c, b); #endif __syncthreads(); diff --git a/hipamd/tests/src/runtimeApi/stream/hipNullStream.cpp b/hipamd/tests/src/runtimeApi/stream/hipNullStream.cpp index 9e6b7d3ff2..26a15cfdb3 100644 --- a/hipamd/tests/src/runtimeApi/stream/hipNullStream.cpp +++ b/hipamd/tests/src/runtimeApi/stream/hipNullStream.cpp @@ -26,11 +26,15 @@ THE SOFTWARE. #include "hip/hip_runtime.h" #include "test_common.h" + +#include #include + unsigned p_streams = 16; int p_repeat = 10; int p_db = 0; +using namespace std; template __global__ void vectorADDRepeat(hipLaunchParm lp, const T* A_d, const T* B_d, T* C_d, size_t NELEM,