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..61fd9bdbe7 --- /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; + + //! 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); 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_fp16_math_fwd.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h index fccbcfbfdc..83fd1349b1 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h +++ b/projects/clr/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); 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 d4207c56fe..3db06bb15e 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 ; @@ -110,9 +110,8 @@ extern int HIP_TRACE_API; #include #include #if __HCC__ -#include -#endif // __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/include/hip/hcc_detail/math_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h index b77ebc8f76..b12e7aca89 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h +++ b/projects/clr/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 + +#include #include #include #include #include -#include - -#include - -#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); } 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/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..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 5849e4e6c6..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 9ed016c7ad..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 4db3b2fea8..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; diff --git a/projects/clr/hipamd/src/device_util.cpp b/projects/clr/hipamd/src/device_util.cpp index 7fa77dc5fe..c86e52617b 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/hipIntegerIntrinsics.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp index f75e406e98..b8fe1d27c4 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/hipSinglePrecisionIntrinsics.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp index cad06fe4ec..f3d2a36931 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp +++ b/projects/clr/hipamd/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp @@ -27,6 +27,7 @@ THE SOFTWARE. #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 61876c33dc..3519d597b3 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hip_trig.cpp +++ b/projects/clr/hipamd/tests/src/deviceLib/hip_trig.cpp @@ -28,10 +28,11 @@ THE SOFTWARE. #include #include -#include #include "test_common.h" #include +#include + #define HIP_ASSERT(x) (assert((x) == hipSuccess)) #define LEN 512 diff --git a/projects/clr/hipamd/tests/src/kernel/hipLanguageExtensions.cpp b/projects/clr/hipamd/tests/src/kernel/hipLanguageExtensions.cpp index 63655e5092..f40014d399 100644 --- a/projects/clr/hipamd/tests/src/kernel/hipLanguageExtensions.cpp +++ b/projects/clr/hipamd/tests/src/kernel/hipLanguageExtensions.cpp @@ -28,6 +28,7 @@ THE SOFTWARE. */ #include "hip/hip_runtime.h" + #include #ifdef __HCC__ diff --git a/projects/clr/hipamd/tests/src/runtimeApi/stream/hipNullStream.cpp b/projects/clr/hipamd/tests/src/runtimeApi/stream/hipNullStream.cpp index ad16455615..bf4699ddb6 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(const T* A_d, const T* B_d, T* C_d, size_t NELEM,