Align with HC Next.

[ROCm/clr commit: 26c26f63cb]
2018-09-17 11:50:29 +03:00
parent 191a752474
commit f5d204450c
15 changed files with 149 additions and 51 deletions
@@ -0,0 +1,69 @@
+#pragma once
+
+#include <stdint.h>
+
+#include <hc_defines.h>
+
+#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);
@@ -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;
+}
+
@@ -71,7 +71,7 @@ THE SOFTWARE.
 //---
 // Remainder of this file only compiles with HCC
 #if defined __HCC__
-#include <grid_launch.h>
+#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 <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>
-#endif // __HCC__

 // TODO-HCC remove old definitions ; ~1602 hcc supports __HCC_ACCELERATOR__ define.
 #if defined(__KALMAR_ACCELERATOR__) && !defined(__HCC_ACCELERATOR__)
@@ -26,17 +26,12 @@ THE SOFTWARE.

 #include <hip/hcc_detail/host_defines.h>

+#include <algorithm>
 #include <assert.h>
 #include <limits.h>
 #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")

@@ -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); }
@@ -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))
@@ -21,10 +21,9 @@ THE SOFTWARE.
 */

 #include <hc.hpp>
-#include <grid_launch.h>
-#include <hc_math.hpp>
 #include "device_util.h"
 #include "hip/hcc_detail/device_functions.h"
+#include "hip/hcc_detail/grid_launch.h"
 #include "hip/hip_runtime.h"
 #include <atomic>

@@ -28,6 +28,7 @@ THE SOFTWARE.

 #include <hip/hip_runtime.h>
 #include <hip/device_functions.h>
+#include <hip/math_functions.h>
 #include "test_common.h"

 #pragma GCC diagnostic ignored "-Wall"
@@ -27,6 +27,7 @@ THE SOFTWARE.
 */

 #include "test_common.h"
+#include <hip/math_functions.h>
 #include <hip/device_functions.h>

 #define LEN 512
@@ -31,6 +31,10 @@ THE SOFTWARE.
 #include <hip/device_functions.h>
 #include "test_common.h"

+#include <algorithm>
+
+using namespace std;
+
 #pragma GCC diagnostic ignored "-Wall"
 #pragma clang diagnostic ignored "-Wunused-variable"

@@ -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__
@@ -215,6 +215,7 @@ template<
    typename T, 
    typename enable_if<
        is_same<T, int>{} || is_same<T, unsigned int>{}>::type* = nullptr>
+__device__
 void testKernelSub(T* g_odata) {
    // Atomic subtraction (final should be 0)
    atomicSub(&g_odata[1], 10);
@@ -27,6 +27,8 @@ THE SOFTWARE.

 #include <hip/hip_runtime.h>
 #include <hip/device_functions.h>
+#include <hip/math_functions.h>
+
 #include "test_common.h"

 #pragma GCC diagnostic ignored "-Wall"
@@ -28,9 +28,11 @@ THE SOFTWARE.

 #include <hip/hip_runtime_api.h>
 #include <hip/hip_runtime.h>
-#include <iostream>
 #include "test_common.h"
 #include <hip/device_functions.h>
+#include <hip/math_functions.h>
+
+#include <iostream>

 #define LEN 512
 #define SIZE LEN << 2
@@ -31,7 +31,7 @@ THE SOFTWARE.
 #include <test_common.h>

 #ifdef __HCC__
-#include <amp.h>
+    #include <hc.hpp>
 #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();
@@ -26,11 +26,15 @@ THE SOFTWARE.

 #include "hip/hip_runtime.h"
 #include "test_common.h"
+
+#include <algorithm>
 #include <vector>
+
 unsigned p_streams = 16;
 int p_repeat = 10;
 int p_db = 0;

+using namespace std;

 template <typename T>
 __global__ void vectorADDRepeat(hipLaunchParm lp, const T* A_d, const T* B_d, T* C_d, size_t NELEM,