From f5d204450c89cbdd57aeca7bfe40c6f7ff0cad07 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Mon, 17 Sep 2018 11:50:29 +0300 Subject: [PATCH 1/8] Align with HC Next. [ROCm/clr commit: 26c26f63cb12405b92947ce4989dc380de6d7cff] --- .../include/hip/hcc_detail/grid_launch.h | 69 +++++++++++++++++++ .../include/hip/hcc_detail/grid_launch.hpp | 50 ++++++++++++++ .../include/hip/hcc_detail/hip_runtime.h | 5 +- .../include/hip/hcc_detail/math_functions.h | 23 +++---- .../hipamd/include/hip/hcc_detail/math_fwd.h | 24 ------- projects/clr/hipamd/src/device_util.cpp | 3 +- .../hipDoublePrecisionIntrinsics.cpp | 1 + .../tests/src/deviceLib/hipFloatMath.cpp | 1 + .../src/deviceLib/hipIntegerIntrinsics.cpp | 4 ++ .../tests/src/deviceLib/hipMathFunctions.cpp | 1 + .../src/deviceLib/hipSimpleAtomicsTest.cpp | 1 + .../hipSinglePrecisionIntrinsics.cpp | 2 + .../hipamd/tests/src/deviceLib/hip_trig.cpp | 4 +- .../src/kernel/hipLanguageExtensions.cpp | 8 +-- .../src/runtimeApi/stream/hipNullStream.cpp | 4 ++ 15 files changed, 149 insertions(+), 51 deletions(-) create mode 100644 projects/clr/hipamd/include/hip/hcc_detail/grid_launch.h create mode 100644 projects/clr/hipamd/include/hip/hcc_detail/grid_launch.hpp diff --git a/projects/clr/hipamd/include/hip/hcc_detail/grid_launch.h b/projects/clr/hipamd/include/hip/hcc_detail/grid_launch.h new file mode 100644 index 0000000000..f91d23341a --- /dev/null +++ b/projects/clr/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/projects/clr/hipamd/include/hip/hcc_detail/grid_launch.hpp b/projects/clr/hipamd/include/hip/hcc_detail/grid_launch.hpp new file mode 100644 index 0000000000..04ce7e0366 --- /dev/null +++ b/projects/clr/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/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h index 7a6aba5ab3..9759a9bdf8 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/projects/clr/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/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h index 6c8510fcbb..c7a4d0e88f 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h +++ b/projects/clr/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/projects/clr/hipamd/include/hip/hcc_detail/math_fwd.h b/projects/clr/hipamd/include/hip/hcc_detail/math_fwd.h index 102714e117..404c2f81d5 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/math_fwd.h +++ b/projects/clr/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/projects/clr/hipamd/src/device_util.cpp b/projects/clr/hipamd/src/device_util.cpp index 87fbe0fcbc..2e78dea009 100644 --- a/projects/clr/hipamd/src/device_util.cpp +++ b/projects/clr/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/projects/clr/hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp index 76b7a5de7a..413f67d012 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp +++ b/projects/clr/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/projects/clr/hipamd/tests/src/deviceLib/hipFloatMath.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipFloatMath.cpp index 157ea046ec..403a77f703 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipFloatMath.cpp +++ b/projects/clr/hipamd/tests/src/deviceLib/hipFloatMath.cpp @@ -27,6 +27,7 @@ THE SOFTWARE. */ #include "test_common.h" +#include #include #define LEN 512 diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp index 1e6a645b1a..d62a591257 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp +++ b/projects/clr/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/projects/clr/hipamd/tests/src/deviceLib/hipMathFunctions.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipMathFunctions.cpp index f89bdae149..3acafddcc0 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipMathFunctions.cpp +++ b/projects/clr/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/projects/clr/hipamd/tests/src/deviceLib/hipSimpleAtomicsTest.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipSimpleAtomicsTest.cpp index 129d7c1aa8..4890ed0f97 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipSimpleAtomicsTest.cpp +++ b/projects/clr/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/projects/clr/hipamd/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp index dab153ed5e..c0d21aade8 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp +++ b/projects/clr/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/projects/clr/hipamd/tests/src/deviceLib/hip_trig.cpp b/projects/clr/hipamd/tests/src/deviceLib/hip_trig.cpp index 3f200fcf8f..e9bc1dcf13 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hip_trig.cpp +++ b/projects/clr/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/projects/clr/hipamd/tests/src/kernel/hipLanguageExtensions.cpp b/projects/clr/hipamd/tests/src/kernel/hipLanguageExtensions.cpp index 36b79965ca..2677d4da8e 100644 --- a/projects/clr/hipamd/tests/src/kernel/hipLanguageExtensions.cpp +++ b/projects/clr/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/projects/clr/hipamd/tests/src/runtimeApi/stream/hipNullStream.cpp b/projects/clr/hipamd/tests/src/runtimeApi/stream/hipNullStream.cpp index 9e6b7d3ff2..26a15cfdb3 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/stream/hipNullStream.cpp +++ b/projects/clr/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, From 16b5a7c63eb6b47a63be533facb8e26d432a4593 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Wed, 10 Oct 2018 11:56:54 +0100 Subject: [PATCH 2/8] Use HIP math functions, stop using hipLaunchParm. [ROCm/clr commit: 031c1d574f18619d22948ac9c1046bbefcb963e1] --- .../clr/hipamd/tests/src/kernel/hipLanguageExtensions.cpp | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/projects/clr/hipamd/tests/src/kernel/hipLanguageExtensions.cpp b/projects/clr/hipamd/tests/src/kernel/hipLanguageExtensions.cpp index 2677d4da8e..d9b35a15ec 100644 --- a/projects/clr/hipamd/tests/src/kernel/hipLanguageExtensions.cpp +++ b/projects/clr/hipamd/tests/src/kernel/hipLanguageExtensions.cpp @@ -28,6 +28,8 @@ THE SOFTWARE. */ #include "hip/hip_runtime.h" +#include "hip/math_functions.h" + #include #ifdef __HCC__ @@ -53,7 +55,7 @@ __device__ __forceinline__ int sum1_forceinline(int a) { return a + 1; }; __device__ __host__ float PlusOne(float x) { return x + 1.0; } -__global__ void MyKernel(const hipLaunchParm lp, const float* a, const float* b, float* c, +__global__ void MyKernel(const float* a, const float* b, float* c, unsigned N) { // KERNELBEGIN; @@ -71,12 +73,12 @@ void callMyKernel() { const unsigned blockSize = 256; unsigned N = blockSize; - hipLaunchKernel(MyKernel, dim3(N / blockSize), dim3(blockSize), 0, 0, a, b, c, N); + hipLaunchKernelGGL(MyKernel, dim3(N / blockSize), dim3(blockSize), 0, 0, a, b, c, N); } template -__global__ void vectorADD(const hipLaunchParm lp, T __restrict__* A_d, T* B_d, T* C_d, size_t N) { +__global__ void vectorADD(T __restrict__* A_d, T* B_d, T* C_d, size_t N) { // KERNELBEGIN; #ifdef NOT_YET int a = __shfl_up(x, 1); From 21997797fff030c5880792a28ce6a5df14bdf842 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Thu, 11 Oct 2018 00:03:01 +0100 Subject: [PATCH 3/8] Address Aaron's comments [ROCm/clr commit: b2b482c78ed01be799c7e72da67ddcbb03b49394] --- projects/clr/hipamd/include/hip/hcc_detail/grid_launch.h | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/grid_launch.h b/projects/clr/hipamd/include/hip/hcc_detail/grid_launch.h index f91d23341a..61fd9bdbe7 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/grid_launch.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/grid_launch.h @@ -34,11 +34,11 @@ typedef struct grid_launch_parm gl_dim3 grid_dim; //! Group dimensions - gl_dim3 group_dim;; + 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;; + unsigned int dynamic_group_mem_bytes; //! Control setting of barrier bit on per-packet basis: //! See gl_barrier_bit description. @@ -56,7 +56,7 @@ typedef struct grid_launch_parm //! 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. + //! Pointer 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. From 0e5f21c96b0337d4d3c93e1f9de033e3367dd458 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Thu, 11 Oct 2018 00:08:09 +0100 Subject: [PATCH 4/8] Trailing whitespace is still not the final frontier. [ROCm/clr commit: 7874ec8aafe714ce7ac69789f3e7abf424e85d97] --- projects/clr/hipamd/tests/src/deviceLib/hipTestFMA.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipTestFMA.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipTestFMA.cpp index 6350999084..44b96d518d 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipTestFMA.cpp +++ b/projects/clr/hipamd/tests/src/deviceLib/hipTestFMA.cpp @@ -62,7 +62,7 @@ THE SOFTWARE. Check(fma(d, (double)f, (unsigned char)3) == doubleResult); Check(fma(d, (double)f, (short)3) == doubleResult); Check(fma(d, (double)f, (unsigned short)3) == doubleResult); - Check(fma(d, (double)f, (int)3) == doubleResult); + Check(fma(d, (double)f, (int)3) == doubleResult); Check(fma(d, (double)f, (unsigned int)3) == doubleResult); Check(fma(d, (double)f, (long)3) == doubleResult); Check(fma(d, (double)f, (unsigned long)3) == doubleResult); From f3495964ea7d951304f2acda25582c6b78d0a0ec Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Thu, 11 Oct 2018 00:21:41 +0100 Subject: [PATCH 5/8] Minimal should mean minimal. [ROCm/clr commit: 217fe345cac559b48b403f62f9c114ede3f8e7be] --- projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h | 4 +++- .../tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp | 1 - projects/clr/hipamd/tests/src/deviceLib/hipFloatMath.cpp | 1 - projects/clr/hipamd/tests/src/deviceLib/hipMathFunctions.cpp | 1 - .../tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp | 1 - projects/clr/hipamd/tests/src/deviceLib/hipTestFMA.cpp | 1 - projects/clr/hipamd/tests/src/deviceLib/hip_trig.cpp | 1 - .../clr/hipamd/tests/src/kernel/hipLanguageExtensions.cpp | 1 - 8 files changed, 3 insertions(+), 8 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h index cf4125be31..3db06bb15e 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -109,7 +109,9 @@ extern int HIP_TRACE_API; #include #include #include - +#if __HCC__ + #include +#endif // TODO-HCC remove old definitions ; ~1602 hcc supports __HCC_ACCELERATOR__ define. #if defined(__KALMAR_ACCELERATOR__) && !defined(__HCC_ACCELERATOR__) #define __HCC_ACCELERATOR__ __KALMAR_ACCELERATOR__ diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp index 413f67d012..76b7a5de7a 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp +++ b/projects/clr/hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp @@ -28,7 +28,6 @@ THE SOFTWARE. #include #include -#include #include "test_common.h" #pragma GCC diagnostic ignored "-Wall" diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipFloatMath.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipFloatMath.cpp index fb7e4451dc..d363efc271 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipFloatMath.cpp +++ b/projects/clr/hipamd/tests/src/deviceLib/hipFloatMath.cpp @@ -27,7 +27,6 @@ THE SOFTWARE. */ #include "test_common.h" -#include #include #define LEN 512 diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipMathFunctions.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipMathFunctions.cpp index 3acafddcc0..f89bdae149 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipMathFunctions.cpp +++ b/projects/clr/hipamd/tests/src/deviceLib/hipMathFunctions.cpp @@ -27,7 +27,6 @@ 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/projects/clr/hipamd/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp index c0d21aade8..54f8118423 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp +++ b/projects/clr/hipamd/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp @@ -27,7 +27,6 @@ THE SOFTWARE. #include #include -#include #include "test_common.h" diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipTestFMA.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipTestFMA.cpp index 44b96d518d..2771fac585 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipTestFMA.cpp +++ b/projects/clr/hipamd/tests/src/deviceLib/hipTestFMA.cpp @@ -26,7 +26,6 @@ THE SOFTWARE. #include "test_common.h" #include #include -#include #include #define HIP_ASSERT(status) assert(status == hipSuccess) diff --git a/projects/clr/hipamd/tests/src/deviceLib/hip_trig.cpp b/projects/clr/hipamd/tests/src/deviceLib/hip_trig.cpp index 59eaeacadb..60f54a7a3d 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hip_trig.cpp +++ b/projects/clr/hipamd/tests/src/deviceLib/hip_trig.cpp @@ -30,7 +30,6 @@ THE SOFTWARE. #include #include "test_common.h" #include -#include #include diff --git a/projects/clr/hipamd/tests/src/kernel/hipLanguageExtensions.cpp b/projects/clr/hipamd/tests/src/kernel/hipLanguageExtensions.cpp index d9b35a15ec..e9c8629662 100644 --- a/projects/clr/hipamd/tests/src/kernel/hipLanguageExtensions.cpp +++ b/projects/clr/hipamd/tests/src/kernel/hipLanguageExtensions.cpp @@ -28,7 +28,6 @@ THE SOFTWARE. */ #include "hip/hip_runtime.h" -#include "hip/math_functions.h" #include From 2577b11365e2adddab5584e158f227c7248089a2 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Sat, 13 Oct 2018 23:28:22 +0100 Subject: [PATCH 6/8] Update samples. [ROCm/clr commit: f480a134f491ab81fa5afa268af556e2bff12457] --- .../hcc_dialects/vadd_amp_arrayview.cpp | 74 ------------------- .../0_Intro/hcc_dialects/vadd_hc_am.cpp | 4 +- .../0_Intro/hcc_dialects/vadd_hc_array.cpp | 2 +- .../0_Intro/hcc_dialects/vadd_hc_array.hc | 2 +- .../hcc_dialects/vadd_hc_arrayview.cpp | 2 +- 5 files changed, 5 insertions(+), 79 deletions(-) delete mode 100644 projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_amp_arrayview.cpp diff --git a/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_amp_arrayview.cpp b/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_amp_arrayview.cpp deleted file mode 100644 index bb0a4157f8..0000000000 --- a/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_amp_arrayview.cpp +++ /dev/null @@ -1,74 +0,0 @@ -/* -Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. - -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: - -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. - -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ - -// Simple test showing how to use C++AMP syntax with array_view. -// The code uses AMP's array_view class, which provides automatic data synchronization -// of data between the host and the accelerator. As noted below, the HCC runtime -// will automatically copy data to and from the host, without the user needing -// to manually perform such copies. This is an excellent mode for developers -// new to GPU programming and matches the memory models provided by recent systems where -// CPU and GPU share the same memory pool. Advanced programmers may prefer -// more explicit control over the data movement - shown in the other vadd_hc_array and -// vadd_hc_am examples. -// This example shows the similarity between C++AMP and and HC for simple cases where -// implicit data transfer is used - really the only difference is the namespace. -// Other examples show some of the more advanced controls. - -#include - -int main(int argc, char* argv[]) { - int sizeElements = 1000000; - bool pass = true; - - // Allocate auto-managed host/device views of data: - concurrency::array_view A(sizeElements); - concurrency::array_view B(sizeElements); - concurrency::array_view C(sizeElements); - - // Initialize host data - for (int i = 0; i < sizeElements; i++) { - A[i] = 1.618f * i; - B[i] = 3.142f * i; - } - C.discard_data(); // tell runtime not to copy CPU host data. - - - // Launch kernel onto default accelerator - // The HCC runtime will ensure that A and B are available on the accelerator before launching - // the kernel. - concurrency::parallel_for_each(concurrency::extent<1>(sizeElements), - [=](concurrency::index<1> idx) restrict(amp) { - int i = idx[0]; - C[i] = A[i] + B[i]; - }); - - for (int i = 0; i < sizeElements; i++) { - float ref = 1.618f * i + 3.142f * i; - // Because C is an array_view, the HCC runtime will copy C back to host at first access - // here: - if (C[i] != ref) { - printf("error:%d computed=%6.2f, reference=%6.2f\n", i, C[i], ref); - pass = false; - } - }; - if (pass) printf("PASSED!\n"); -} diff --git a/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp b/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp index 65b38df099..9cfad779f9 100644 --- a/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp +++ b/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp @@ -26,8 +26,8 @@ THE SOFTWARE. // which can only be used on the device. The programmer has full control // over when data is copied. -#include -#include +#include +#include int main(int argc, char* argv[]) { int sizeElements = 1000000; diff --git a/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_array.cpp b/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_array.cpp index 5849e4e6c6..aff64528aa 100644 --- a/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_array.cpp +++ b/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_array.cpp @@ -27,7 +27,7 @@ THE SOFTWARE. // automatic data management capabilities - instead the programmer // takes the reins and controls when copies are executed. -#include +#include int main(int argc, char* argv[]) { int sizeElements = 1000000; diff --git a/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_array.hc b/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_array.hc index 9ed016c7ad..cd3d32fd4c 100644 --- a/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_array.hc +++ b/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_array.hc @@ -20,7 +20,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include +#include int main(int argc, char *argv[]) { diff --git a/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_arrayview.cpp b/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_arrayview.cpp index 4db3b2fea8..2d2089abc7 100644 --- a/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_arrayview.cpp +++ b/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_arrayview.cpp @@ -33,7 +33,7 @@ THE SOFTWARE. // implicit data transfer is used - really the only difference is the namespace. // Other examples show some of the more advanced controls. -#include +#include int main(int argc, char* argv[]) { int sizeElements = 1000000; From fe7e09d222ae5551b9e8f03f2fa039f10232a7af Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Thu, 18 Oct 2018 12:29:27 +0100 Subject: [PATCH 7/8] Guard new includes. [ROCm/clr commit: 59533b84e1b437cf73f3a38590d272d11e423343] --- .../hipamd/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp | 9 +++++++-- .../samples/0_Intro/hcc_dialects/vadd_hc_array.cpp | 6 +++++- .../hipamd/samples/0_Intro/hcc_dialects/vadd_hc_array.hc | 6 +++++- .../samples/0_Intro/hcc_dialects/vadd_hc_arrayview.cpp | 6 +++++- 4 files changed, 22 insertions(+), 5 deletions(-) diff --git a/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp b/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp index 9cfad779f9..47fd46cc40 100644 --- a/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp +++ b/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp @@ -26,8 +26,13 @@ THE SOFTWARE. // which can only be used on the device. The programmer has full control // over when data is copied. -#include -#include +#if defined(HC_NEXT) + #include + #include +#else + #include + #include +#endif int main(int argc, char* argv[]) { int sizeElements = 1000000; diff --git a/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_array.cpp b/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_array.cpp index aff64528aa..3599eb07db 100644 --- a/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_array.cpp +++ b/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_array.cpp @@ -27,7 +27,11 @@ THE SOFTWARE. // automatic data management capabilities - instead the programmer // takes the reins and controls when copies are executed. -#include +#if defined(HC_NEXT) + #include +#else + #include +#endif int main(int argc, char* argv[]) { int sizeElements = 1000000; diff --git a/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_array.hc b/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_array.hc index cd3d32fd4c..605d11f999 100644 --- a/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_array.hc +++ b/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_array.hc @@ -20,7 +20,11 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include +#if defined(HC_NEXT) + #include +#else + #include +#endif int main(int argc, char *argv[]) { diff --git a/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_arrayview.cpp b/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_arrayview.cpp index 2d2089abc7..46ac8427a7 100644 --- a/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_arrayview.cpp +++ b/projects/clr/hipamd/samples/0_Intro/hcc_dialects/vadd_hc_arrayview.cpp @@ -33,7 +33,11 @@ THE SOFTWARE. // implicit data transfer is used - really the only difference is the namespace. // Other examples show some of the more advanced controls. -#include +#if defined(HC_NEXT) + #include +#else + #include +#endif int main(int argc, char* argv[]) { int sizeElements = 1000000; From 35e9dfc593c709bc7df06e356e70717ed02ed76e Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Thu, 18 Oct 2018 15:33:46 +0100 Subject: [PATCH 8/8] Dumb workaround is still needed, so add it back. [ROCm/clr commit: 3678063598540690576b381c869dbfd1007296bd] --- projects/clr/hipamd/include/hip/hcc_detail/math_functions.h | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h index faf3804719..4a68a938af 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h @@ -33,6 +33,12 @@ THE SOFTWARE. #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")