Merge pull request #705 from ROCm-Developer-Tools/feature_minimal_changes_for_hc_next

Feature minimal changes for hc next

hipamd/include/hip/hcc_detail/grid_launch.h

@@ -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;
+
+  //! 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.
+  hc::completion_future *cf;
+
+  grid_launch_parm() = default;
+} grid_launch_parm;
+
+
+extern void init_grid_launch(grid_launch_parm *gl);

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;
+}
+

hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h

@@ -26,6 +26,8 @@ THE SOFTWARE.
 // Half Math Functions
 // */
 
+#include "host_defines.h"
+
 extern "C"
 {
     __device__ __attribute__((const)) _Float16 __ocml_ceil_f16(_Float16);

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 <grid_launch.h>
+#include "grid_launch.h"
 #include "hc_printf.hpp"
 // TODO-HCC-GL - change this to typedef.
 // typedef grid_launch_parm hipLaunchParm ;
@@ -110,9 +110,8 @@ extern int HIP_TRACE_API;
 #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__
-
+    #include <hip/hcc_detail/math_functions.h>
+#endif
 // TODO-HCC remove old definitions ; ~1602 hcc supports __HCC_ACCELERATOR__ define.
 #if defined(__KALMAR_ACCELERATOR__) && !defined(__HCC_ACCELERATOR__)
 #define __HCC_ACCELERATOR__ __KALMAR_ACCELERATOR__

hipamd/include/hip/hcc_detail/math_functions.h

@@ -22,16 +22,16 @@ THE SOFTWARE.
 
 #pragma once
 
+#include "hip_fp16_math_fwd.h"
+#include "math_fwd.h"
+
+#include <hip/hcc_detail/host_defines.h>
+
+#include <algorithm>
 #include <assert.h>
 #include <limits.h>
 #include <limits>
 #include <stdint.h>
-#include <algorithm>
-
-#include <hip/hcc_detail/host_defines.h>
-
-#include "hip_fp16_math_fwd.h"
-#include "math_fwd.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.
@@ -594,16 +594,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); }
@@ -1084,16 +1084,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); }

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))