From 34ee235d0627419ccc8089122e8d5f671f00d4bb Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Thu, 28 Jan 2016 09:51:11 -0600 Subject: [PATCH 01/16] Fix typo in hipStreamWaitEvent. Fixes#9 [ROCm/hip commit: 1e2fb8a1abdb554aa3c6462178e0bf947ea476f2] --- projects/hip/bin/hipify | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/bin/hipify b/projects/hip/bin/hipify index 5cab931998..3eecc12abb 100755 --- a/projects/hip/bin/hipify +++ b/projects/hip/bin/hipify @@ -325,7 +325,7 @@ while (@ARGV) { $ft{'stream'} += s/\bcudaStreamCreate\b/hipStreamCreate/g; $ft{'stream'} += s/\bcudaStreamCreateWithFlags\b/hipStreamCreateWithFlags/g; $ft{'stream'} += s/\bcudaStreamDestroy\b/hipStreamDestroy/g; - $ft{'stream'} += s/\bcudaStreamWaitEvent\b/hipStreamWaitEven/g; + $ft{'stream'} += s/\bcudaStreamWaitEvent\b/hipStreamWaitEvent/g; $ft{'stream'} += s/\bcudaStreamSynchronize\b/hipStreamSynchronize/g; $ft{'stream'} += s/\bcudaStreamDefault\b/hipStreamDefault/g; $ft{'stream'} += s/\bcudaStreamNonBlocking\b/hipStreamNonBlocking/g; From b37fdaff32abf020e5657677d1fa847193a9bef4 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Paul=20J=C3=A4hne?= Date: Sun, 31 Jan 2016 13:29:03 +0100 Subject: [PATCH 02/16] Update hip_faq.md changed Cuda to CUDA for consistency changed NVIDIA to Nvidia for consistency corrected apostrophes corrected section "What hardware does HIP support?" [ROCm/hip commit: c53031809cd88410d686b996a6b67be4f3740f63] --- projects/hip/docs/markdown/hip_faq.md | 36 +++++++++++++-------------- 1 file changed, 18 insertions(+), 18 deletions(-) diff --git a/projects/hip/docs/markdown/hip_faq.md b/projects/hip/docs/markdown/hip_faq.md index 969982f21e..a6a8bb1661 100644 --- a/projects/hip/docs/markdown/hip_faq.md +++ b/projects/hip/docs/markdown/hip_faq.md @@ -8,11 +8,11 @@ HIP provides the following: - Streams (hipStreamCreate(), etc.)---under development - Events (hipEventRecord(), hipEventElapsedTime(), etc.) - Kernel launching (hipLaunchKernel is a standard C/C++ function that replaces <<< >>>) -- Cuda-style kernel indexing +- CUDA-style kernel indexing - Device-side math built-ins - Error reporting (hipGetLastError(), hipGetErrorString()) -The HIP documentation describes each API and its limitations, if any, compared with the equivalent Cuda API. +The HIP documentation describes each API and its limitations, if any, compared with the equivalent CUDA API. ### What is not supported? #### Run-time features: @@ -20,8 +20,8 @@ The HIP documentation describes each API and its limitations, if any, compared w - Dynamic parallelism - Managed memory - Graphics interoperation with OpenGL or Direct3D -- Cuda array, mipmappedArray and pitched memory -- Cuda Driver API +- CUDA array, mipmappedArray and pitched memory +- CUDA Driver API #### Kernel language features: - Device-side dynamic memory allocations (malloc, free, new, delete) @@ -34,35 +34,35 @@ The HIP documentation describes each API and its limitations, if any, compared w Both AMD and Nvidia support OpenCL 1.2 on their devices, so developers can write portable code. HIP offers several benefits over OpenCL: - Developers can code in C++ as well as mix host and device C++ code in their source files. HIP C++ code can use templates, lambdas, classes and so on. -- The HIP API is less verbose than OpenCL and is familiar to Cuda developers. -- Because both Cuda and HIP are C++ languages, porting from Cuda to HIP is significantly easier than porting from Cuda to OpenCL. +- The HIP API is less verbose than OpenCL and is familiar to CUDA developers. +- Because both CUDA and HIP are C++ languages, porting from CUDA to HIP is significantly easier than porting from CUDA to OpenCL. - HIP uses the best available development tools on each platform: on Nvidia GPUs, HIP code compiles using NVCC and can employ the nSight profiler and debugger (unlike OpenCL on Nvidia GPUs). - HIP provides pointers and host-side pointer arithmetic. - HIP provides device-level control over memory allocation and placement. - HIP offers an offline compilation model. ### What hardware does HIP support? -- For AMD platforms, HIP runs on the same hardware that the HCC "hc" mode supports---specifically AMD Kaveri, Carrizo and Fiji. For Nvidia platforms, it should run on a device that uses the Cuda SDK 5.0 or newer. We have tested the Nvidia Titan and K40. -- For NVIDIA platform, HIP should run on an device which runs the CUDA SDK 5.0 or newer. We have tested nvidia Titan and K40. +- For AMD platforms, HIP runs on the same hardware that the HCC "hc" mode supports - specifically AMD Kaveri, Carrizo and Fiji. +- For Nvidia platforms, HIP should run on a device which runs the CUDA SDK 5.0 or newer. We have tested the Nvidia Titan and K40. ### Does Hipify automatically convert all source code? Typically, Hipify can automatically convert almost all run-time code, and the coordinate indexing device code. -Most device code needs no additional conversion, since HIP and Cuda have similar names for math and built-in functions. +Most device code needs no additional conversion, since HIP and CUDA have similar names for math and built-in functions. HIP currently requires manual addition of one more arguments to the kernel so that the host can communicate the execution configuration to the device. -Additional porting may be required to deal with architecture feature queries or with Cuda capabilities that HIP doesn’t support. +Additional porting may be required to deal with architecture feature queries or with CUDA capabilities that HIP doesn't support. Developers should always expect to perform some platform-specific tuning and optimization. ### What is NVCC? -NVCC is Nvidia's compiler driver for compiling "Cuda C++" code into PTX or device code for Nvidia GPUs. It’s a closed-source binary product that comes with Cuda SDKs. +NVCC is Nvidia's compiler driver for compiling "CUDA C++" code into PTX or device code for Nvidia GPUs. It's a closed-source binary product that comes with CUDA SDKs. ### What is HCC? HCC is AMD's compiler driver which compiles "heterogenous C++" code into HSAIL or GCN device code for AMD GPUs. HCC is an open-source compiler based on recent versions of CLANG/LLVM. -### Why use HIP rather than supporting Cuda run time directly? +### Why use HIP rather than supporting CUDA run time directly? While HIP is a strong subset of the CUDA, it is a subset. The HIP layer allows that subset to be clearly defined and documented. -Developers who code to the HIP API can be assured there code will remain portable across NVIDIA and AMD platforms. +Developers who code to the HIP API can be assured there code will remain portable across Nvidia and AMD platforms. -### Can I develop HIP code on an NVIDIA CUDA platform? +### Can I develop HIP code on an Nvidia CUDA platform? Yes! HIP's CUDA path only exposes the APIs and functionality that work on both NVCC and HCC back-ends. "Extra" APIs, parameters, and features which exist in CUDA but not in HCC will typically result in compile-time or run-time errors. Developers need to use the HIP API for most accelerator code, and bracket any CUDA-specific code with appropriate ifdefs. @@ -70,12 +70,12 @@ Developers concerned about portability should of course run on both platforms, a In some cases CUDA has a richer set of modes for some APIs, and some C++ capabilities such as virtual functions - see the HIP @API documentation for more details. ### Can I develop HIP code on an AMD HCC platform? -Yes! HIP's HCC path only exposes the APIs and functions that work on both NVCC and HCC back ends. "Extra" APIs, parameters and features that appear in HCC but not Cuda will typically cause compile- or run-time errors. Developers must use the HIP API for most accelerator code and bracket any HCC-specific code with appropriate ifdefs. Those concerned about portability should, of course, test their code on both platforms and should tune it for performance. Typically, HCC supports a more modern set of C++11/C++14/C++17 features, so HIP developers who want portability should be careful when using advanced C++ features on the hc path. +Yes! HIP's HCC path only exposes the APIs and functions that work on both NVCC and HCC back ends. "Extra" APIs, parameters and features that appear in HCC but not CUDA will typically cause compile- or run-time errors. Developers must use the HIP API for most accelerator code and bracket any HCC-specific code with appropriate ifdefs. Those concerned about portability should, of course, test their code on both platforms and should tune it for performance. Typically, HCC supports a more modern set of C++11/C++14/C++17 features, so HIP developers who want portability should be careful when using advanced C++ features on the hc path. ### Can a HIP binary run on both AMD and Nvidia platforms? -HIP is a source-portable language that can be compiled to run on either the HCC or NVCC platform. HIP tools don’t create a "fat binary" that can run on either platform, however. +HIP is a source-portable language that can be compiled to run on either the HCC or NVCC platform. HIP tools don't create a "fat binary" that can run on either platform, however. ### What's the difference between HIP and hc? -HIP is a portable C++ language that supports a strong subset of the Cuda run-time APIs and device-kernel language. It’s designed to simplify Cuda conversion to portable C++. HIP provides a C-compatible run-time API, C-compatible kernel-launch mechanism, C++ kernel language and pointer-based memory management. +HIP is a portable C++ language that supports a strong subset of the CUDA run-time APIs and device-kernel language. It's designed to simplify CUDA conversion to portable C++. HIP provides a C-compatible run-time API, C-compatible kernel-launch mechanism, C++ kernel language and pointer-based memory management. -A C++ dialect, hc is supported by the AMD HCC compiler. It provides C++ run time, C++ kernel-launch APIs (parallel_for_each), C++ kernel language, and several memory-management options, including pointers, arrays and array_view (with implicit data synchronization). It’s intended to be a leading indicator of the ISO C++ standard. +A C++ dialect, hc is supported by the AMD HCC compiler. It provides C++ run time, C++ kernel-launch APIs (parallel_for_each), C++ kernel language, and several memory-management options, including pointers, arrays and array_view (with implicit data synchronization). It's intended to be a leading indicator of the ISO C++ standard. From 3e13c7dae73389b5b3d7417517344c19f906d134 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 1 Feb 2016 14:24:41 +0530 Subject: [PATCH 03/16] Remove redundant #define __HCC__ in hcc_detail/hip_runtime.h [ROCm/hip commit: e55f3778e0fd98e7125e17b93f7b10d94572bbb6] --- projects/hip/include/hcc_detail/hip_runtime.h | 7 ------- 1 file changed, 7 deletions(-) diff --git a/projects/hip/include/hcc_detail/hip_runtime.h b/projects/hip/include/hcc_detail/hip_runtime.h index 9d5eac2916..72ecc3515c 100644 --- a/projects/hip/include/hcc_detail/hip_runtime.h +++ b/projects/hip/include/hcc_detail/hip_runtime.h @@ -283,7 +283,6 @@ __device__ inline unsigned long long int atomicXor(unsigned long long int* addre return (long long int)hc::atomic_fetch_xor((uint64_t*)address,(uint64_t)val); } -#ifdef __HCC__ #include // integer intrinsic function __poc __clz __ffs __brev __device__ inline unsigned int __popc( unsigned int input) @@ -337,27 +336,21 @@ __device__ inline unsigned long long int __brevll( unsigned long long int input) } // warp vote function __all __any __ballot - __device__ inline int __all( int input) { return hc::__all( input); } - __device__ inline int __any( int input) { return hc::__any( input); } - __device__ inline unsigned long long int __ballot( int input) { return hc::__ballot( input); } -#endif - - #ifdef __HCC_ACCELERATOR__ #include From 01c51ce7349dfe7bf6470b1d748531052d5bbc4e Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 1 Feb 2016 14:26:50 +0530 Subject: [PATCH 04/16] Restrict using namespace hc::precise_math to device only [ROCm/hip commit: 3b19fd578dbd13ed69ea6318719305505c07fdb3] --- projects/hip/include/hcc_detail/hip_runtime.h | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/projects/hip/include/hcc_detail/hip_runtime.h b/projects/hip/include/hcc_detail/hip_runtime.h index 72ecc3515c..9f5c81152a 100644 --- a/projects/hip/include/hcc_detail/hip_runtime.h +++ b/projects/hip/include/hcc_detail/hip_runtime.h @@ -352,10 +352,11 @@ __device__ inline unsigned long long int __ballot( int input) } -#ifdef __HCC_ACCELERATOR__ #include // TODO: Choose whether default is precise math or fast math based on compilation flag. +#ifdef __HCC_ACCELERATOR__ using namespace hc::precise_math; +#endif //TODO: Undo this once min/max functions are supported by hc inline int min(int arg1, int arg2) __attribute((hc,cpu)) { \ @@ -368,9 +369,6 @@ inline int max(int arg1, int arg2) __attribute((hc,cpu)) { \ __device__ inline float __log2f(float x) {return hc::fast_math::log2(x); }; __device__ inline float __powf(float base, float exponent) {return hc::fast_math::powf(base, exponent); }; -#endif - - /** * Kernel launching From 1b4ad3eedfd644daa4142e5729ee1cb5b71b0ba7 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 1 Feb 2016 14:29:50 +0530 Subject: [PATCH 05/16] Add few more single precision intrinsics to hcc_detail/hip_runtime.h [ROCm/hip commit: d2c6125a7c3422a690984c18449764eb6966eb58] --- projects/hip/include/hcc_detail/hip_runtime.h | 14 +++++++++++++- 1 file changed, 13 insertions(+), 1 deletion(-) diff --git a/projects/hip/include/hcc_detail/hip_runtime.h b/projects/hip/include/hcc_detail/hip_runtime.h index 9f5c81152a..7dfa70ed66 100644 --- a/projects/hip/include/hcc_detail/hip_runtime.h +++ b/projects/hip/include/hcc_detail/hip_runtime.h @@ -366,8 +366,20 @@ inline int max(int arg1, int arg2) __attribute((hc,cpu)) { \ //TODO - add a couple fast math operations here, the set here will grow : -__device__ inline float __log2f(float x) {return hc::fast_math::log2(x); }; +__device__ inline float __cosf(float x) {return hc::fast_math::cosf(x); }; +__device__ inline float __expf(float x) {return hc::fast_math::expf(x); }; +__device__ inline float __frsqrt_rn(float x) {return hc::fast_math::rsqrt(x); }; +__device__ inline float __fsqrt_rd(float x) {return hc::fast_math::sqrt(x); }; +__device__ inline float __fsqrt_rn(float x) {return hc::fast_math::sqrt(x); }; +__device__ inline float __fsqrt_ru(float x) {return hc::fast_math::sqrt(x); }; +__device__ inline float __fsqrt_rz(float x) {return hc::fast_math::sqrt(x); }; +__device__ inline float __log10f(float x) {return hc::fast_math::log10f(x); }; +__device__ inline float __log2f(float x) {return hc::fast_math::log2f(x); }; +__device__ inline float __logf(float x) {return hc::fast_math::logf(x); }; __device__ inline float __powf(float base, float exponent) {return hc::fast_math::powf(base, exponent); }; +__device__ inline void __sincosf(float x, float *s, float *c) {return hc::fast_math::sincosf(x, s, c); }; +__device__ inline float __sinf(float x) {return hc::fast_math::sinf(x); }; +__device__ inline float __tanf(float x) {return hc::fast_math::tanf(x); }; /** From bf51551d228e98b221b8a346fbe34870616ccff6 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 1 Feb 2016 14:34:28 +0530 Subject: [PATCH 06/16] Disable testing of unsupported single precision intrinsics [ROCm/hip commit: 276e20d4dbaf5e5a4b229597d50c263e667f1791] --- .../src/hipSinglePrecisionIntrinsics.cpp | 55 ++++++++++--------- 1 file changed, 28 insertions(+), 27 deletions(-) diff --git a/projects/hip/tests/src/hipSinglePrecisionIntrinsics.cpp b/projects/hip/tests/src/hipSinglePrecisionIntrinsics.cpp index 1b8f80ff82..13d2537f51 100644 --- a/projects/hip/tests/src/hipSinglePrecisionIntrinsics.cpp +++ b/projects/hip/tests/src/hipSinglePrecisionIntrinsics.cpp @@ -30,48 +30,49 @@ __device__ void single_precision_intrinsics() float fX, fY; __cosf(0.0f); - __exp10f(0.0f); + //__exp10f(0.0f); __expf(0.0f); - __fadd_rd(0.0f, 1.0f); - __fadd_rn(0.0f, 1.0f); - __fadd_ru(0.0f, 1.0f); - __fadd_rz(0.0f, 1.0f); - __fdiv_rd(4.0f, 2.0f); - __fdiv_rn(4.0f, 2.0f); - __fdiv_ru(4.0f, 2.0f); - __fdiv_rz(4.0f, 2.0f); - __fdividef(4.0f, 2.0f); - __fmaf_rd(1.0f, 2.0f, 3.0f); - __fmaf_rn(1.0f, 2.0f, 3.0f); - __fmaf_ru(1.0f, 2.0f, 3.0f); - __fmaf_rz(1.0f, 2.0f, 3.0f); - __fmul_rd(1.0f, 2.0f); - __fmul_rn(1.0f, 2.0f); - __fmul_ru(1.0f, 2.0f); - __fmul_rz(1.0f, 2.0f); - __frcp_rd(2.0f); - __frcp_rn(2.0f); - __frcp_ru(2.0f); - __frcp_rz(2.0f); + //__fadd_rd(0.0f, 1.0f); + //__fadd_rn(0.0f, 1.0f); + //__fadd_ru(0.0f, 1.0f); + //__fadd_rz(0.0f, 1.0f); + //__fdiv_rd(4.0f, 2.0f); + //__fdiv_rn(4.0f, 2.0f); + //__fdiv_ru(4.0f, 2.0f); + //__fdiv_rz(4.0f, 2.0f); + //__fdividef(4.0f, 2.0f); + //__fmaf_rd(1.0f, 2.0f, 3.0f); + //__fmaf_rn(1.0f, 2.0f, 3.0f); + //__fmaf_ru(1.0f, 2.0f, 3.0f); + //__fmaf_rz(1.0f, 2.0f, 3.0f); + //__fmul_rd(1.0f, 2.0f); + //__fmul_rn(1.0f, 2.0f); + //__fmul_ru(1.0f, 2.0f); + //__fmul_rz(1.0f, 2.0f); + //__frcp_rd(2.0f); + //__frcp_rn(2.0f); + //__frcp_ru(2.0f); + //__frcp_rz(2.0f); __frsqrt_rn(4.0f); __fsqrt_rd(4.0f); __fsqrt_rn(4.0f); __fsqrt_ru(4.0f); __fsqrt_rz(4.0f); - __fsub_rd(2.0f, 1.0f); - __fsub_rn(2.0f, 1.0f); - __fsub_ru(2.0f, 1.0f); - __fsub_rz(2.0f, 1.0f); + //__fsub_rd(2.0f, 1.0f); + //__fsub_rn(2.0f, 1.0f); + //__fsub_ru(2.0f, 1.0f); + //__fsub_rz(2.0f, 1.0f); __log10f(1.0f); __log2f(1.0f); __logf(1.0f); __powf(1.0f, 0.0f); - __saturatef(0.1f); + //__saturatef(0.1f); __sincosf(0.0f, &fX, &fY); __sinf(0.0f); __tanf(0.0f); } + __global__ void compileSinglePrecisionIntrinsics(hipLaunchParm lp, int ignored) { single_precision_intrinsics(); From af51dd0a72e9df20efb71f118d8afc6f18cee10a Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 1 Feb 2016 14:36:50 +0530 Subject: [PATCH 07/16] Split math function tests into several smaller tests [ROCm/hip commit: 4d9e8b8673eae36a4268ce3f39f5acbccc5b3c38] --- projects/hip/tests/src/CMakeLists.txt | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/projects/hip/tests/src/CMakeLists.txt b/projects/hip/tests/src/CMakeLists.txt index 3f2396b315..968b167e7b 100644 --- a/projects/hip/tests/src/CMakeLists.txt +++ b/projects/hip/tests/src/CMakeLists.txt @@ -115,8 +115,10 @@ make_hip_executable (hipLanguageExtensions hipLanguageExtensions.cpp) make_hip_executable (hipGridLaunch hipGridLaunch.cpp) make_hip_executable (hipHcc hipHcc.cpp) make_hip_executable (hipSimpleAtomicsTest hipSimpleAtomicsTest.cpp) -make_hip_executable (hipMathFunctions hipMathFunctions.cpp hipSinglePrecisionMathHost.cpp hipDoublePrecisionMathHost.cpp hipSinglePrecisionMathDevice.cpp hipDoublePrecisionMathDevice.cpp) -target_link_libraries(hipMathFunctions m) +make_hip_executable (hipMathFunctionsHost hipMathFunctions.cpp hipSinglePrecisionMathHost.cpp hipDoublePrecisionMathHost.cpp) +make_hip_executable (hipMathFunctionsDevice hipMathFunctions.cpp hipSinglePrecisionMathDevice.cpp hipDoublePrecisionMathDevice.cpp) +make_hip_executable (hipIntrinsics hipMathFunctions.cpp hipSinglePrecisionIntrinsics.cpp) +target_link_libraries(hipMathFunctionsHost m) make_test(hip_anyall " " ) make_test(hip_popc " " ) From 79a0c836be23617695c16a7fb168af285bb2cc1b Mon Sep 17 00:00:00 2001 From: Jack Chung Date: Mon, 1 Feb 2016 17:42:37 +0800 Subject: [PATCH 08/16] Disable sincosf which has trouble on hcc now. [ROCm/hip commit: 93d1801a30d2cee3e94a6baa86876e7253b40a68] --- projects/hip/tests/src/hipSinglePrecisionIntrinsics.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/tests/src/hipSinglePrecisionIntrinsics.cpp b/projects/hip/tests/src/hipSinglePrecisionIntrinsics.cpp index 13d2537f51..635e4e8e65 100644 --- a/projects/hip/tests/src/hipSinglePrecisionIntrinsics.cpp +++ b/projects/hip/tests/src/hipSinglePrecisionIntrinsics.cpp @@ -67,7 +67,7 @@ __device__ void single_precision_intrinsics() __logf(1.0f); __powf(1.0f, 0.0f); //__saturatef(0.1f); - __sincosf(0.0f, &fX, &fY); + //__sincosf(0.0f, &fX, &fY); __sinf(0.0f); __tanf(0.0f); } From 9ed3ef50fef7af53872a3092a11dccf522afcf6b Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 1 Feb 2016 16:00:45 +0530 Subject: [PATCH 09/16] Add double and integer intrinsics to test [ROCm/hip commit: 861cba6f751a524406f0a53f90e0a43c7715db44] --- projects/hip/include/hcc_detail/hip_runtime.h | 5 +- projects/hip/tests/src/CMakeLists.txt | 2 +- .../src/hipDoublePrecisionIntrinsics.cpp | 48 +++++++++---------- .../hip/tests/src/hipIntegerIntrinsics.cpp | 26 +++++----- 4 files changed, 42 insertions(+), 39 deletions(-) diff --git a/projects/hip/include/hcc_detail/hip_runtime.h b/projects/hip/include/hcc_detail/hip_runtime.h index 7dfa70ed66..688c9e6959 100644 --- a/projects/hip/include/hcc_detail/hip_runtime.h +++ b/projects/hip/include/hcc_detail/hip_runtime.h @@ -380,7 +380,10 @@ __device__ inline float __powf(float base, float exponent) {return hc::fast_math __device__ inline void __sincosf(float x, float *s, float *c) {return hc::fast_math::sincosf(x, s, c); }; __device__ inline float __sinf(float x) {return hc::fast_math::sinf(x); }; __device__ inline float __tanf(float x) {return hc::fast_math::tanf(x); }; - +__device__ inline float __dsqrt_rd(double x) {return hc::fast_math::sqrt(x); }; +__device__ inline float __dsqrt_rn(double x) {return hc::fast_math::sqrt(x); }; +__device__ inline float __dsqrt_ru(double x) {return hc::fast_math::sqrt(x); }; +__device__ inline float __dsqrt_rz(double x) {return hc::fast_math::sqrt(x); }; /** * Kernel launching diff --git a/projects/hip/tests/src/CMakeLists.txt b/projects/hip/tests/src/CMakeLists.txt index 968b167e7b..8da01db50a 100644 --- a/projects/hip/tests/src/CMakeLists.txt +++ b/projects/hip/tests/src/CMakeLists.txt @@ -117,7 +117,7 @@ make_hip_executable (hipHcc hipHcc.cpp) make_hip_executable (hipSimpleAtomicsTest hipSimpleAtomicsTest.cpp) make_hip_executable (hipMathFunctionsHost hipMathFunctions.cpp hipSinglePrecisionMathHost.cpp hipDoublePrecisionMathHost.cpp) make_hip_executable (hipMathFunctionsDevice hipMathFunctions.cpp hipSinglePrecisionMathDevice.cpp hipDoublePrecisionMathDevice.cpp) -make_hip_executable (hipIntrinsics hipMathFunctions.cpp hipSinglePrecisionIntrinsics.cpp) +make_hip_executable (hipIntrinsics hipMathFunctions.cpp hipSinglePrecisionIntrinsics.cpp hipDoublePrecisionIntrinsics.cpp hipIntegerIntrinsics.cpp) target_link_libraries(hipMathFunctionsHost m) make_test(hip_anyall " " ) diff --git a/projects/hip/tests/src/hipDoublePrecisionIntrinsics.cpp b/projects/hip/tests/src/hipDoublePrecisionIntrinsics.cpp index c18a2c6ab7..ad344aedfd 100644 --- a/projects/hip/tests/src/hipDoublePrecisionIntrinsics.cpp +++ b/projects/hip/tests/src/hipDoublePrecisionIntrinsics.cpp @@ -27,34 +27,34 @@ THE SOFTWARE. __device__ void double_precision_intrinsics() { - __dadd_rd(0.0, 1.0); - __dadd_rn(0.0, 1.0); - __dadd_ru(0.0, 1.0); - __dadd_rz(0.0, 1.0); - __ddiv_rd(4.0, 2.0); - __ddiv_rn(4.0, 2.0); - __ddiv_ru(4.0, 2.0); - __ddiv_rz(4.0, 2.0); - __dmul_rd(1.0, 2.0); - __dmul_rn(1.0, 2.0); - __dmul_ru(1.0, 2.0); - __dmul_rz(1.0, 2.0); - __drcp_rd(2.0); - __drcp_rn(2.0); - __drcp_ru(2.0); - __drcp_rz(2.0); + //__dadd_rd(0.0, 1.0); + //__dadd_rn(0.0, 1.0); + //__dadd_ru(0.0, 1.0); + //__dadd_rz(0.0, 1.0); + //__ddiv_rd(0.0, 1.0); + //__ddiv_rn(0.0, 1.0); + //__ddiv_ru(0.0, 1.0); + //__ddiv_rz(0.0, 1.0); + //__dmul_rd(1.0, 2.0); + //__dmul_rn(1.0, 2.0); + //__dmul_ru(1.0, 2.0); + //__dmul_rz(1.0, 2.0); + //__drcp_rd(2.0); + //__drcp_rn(2.0); + //__drcp_ru(2.0); + //__drcp_rz(2.0); __dsqrt_rd(4.0); __dsqrt_rn(4.0); __dsqrt_ru(4.0); __dsqrt_rz(4.0); - __dsub_rd(2.0, 1.0); - __dsub_rn(2.0, 1.0); - __dsub_ru(2.0, 1.0); - __dsub_rz(2.0, 1.0); - __fma_rd(1.0, 2.0, 3.0); - __fma_rn(1.0, 2.0, 3.0); - __fma_ru(1.0, 2.0, 3.0); - __fma_rz(1.0, 2.0, 3.0); + //__dsub_rd(2.0, 1.0); + //__dsub_rn(2.0, 1.0); + //__dsub_ru(2.0, 1.0); + //__dsub_rz(2.0, 1.0); + //__fma_rd(1.0, 2.0, 3.0); + //__fma_rn(1.0, 2.0, 3.0); + //__fma_ru(1.0, 2.0, 3.0); + //__fma_rz(1.0, 2.0, 3.0); } __global__ void compileDoublePrecisionIntrinsics(hipLaunchParm lp, int ignored) diff --git a/projects/hip/tests/src/hipIntegerIntrinsics.cpp b/projects/hip/tests/src/hipIntegerIntrinsics.cpp index 6061ee8b63..ea6bafcb18 100644 --- a/projects/hip/tests/src/hipIntegerIntrinsics.cpp +++ b/projects/hip/tests/src/hipIntegerIntrinsics.cpp @@ -29,25 +29,25 @@ __device__ void integer_intrinsics() { __brev((unsigned int)10); __brevll((unsigned long long)10); - __byte_perm((unsigned int)0, (unsigned int)0, 0); + //__byte_perm((unsigned int)0, (unsigned int)0, 0); __clz((int)10); __clzll((long long)10); __ffs((int)10); __ffsll((long long)10); - __hadd((int)1, (int)3); - __mul24((int)1, (int)2); - __mul64hi((long long)1, (long long)2); - __mulhi((int)1, (int)2); + //__hadd((int)1, (int)3); + //__mul24((int)1, (int)2); + //__mul64hi((long long)1, (long long)2); + //__mulhi((int)1, (int)2); __popc((unsigned int)4); __popcll((unsigned long long)4); - __rhadd((int)1, (int)2); - __sad((int)1, (int)2, 0); - __uhadd((unsigned int)1, (unsigned int)3); - __umul24((unsigned int)1, (unsigned int)2); - __umul64hi((unsigned long long)1, (unsigned long long)2); - __umulhi((unsigned int)1, (unsigned int)2); - __urhadd((unsigned int)1, (unsigned int)2); - __usad((unsigned int)1, (unsigned int)2, 0); + //__rhadd((int)1, (int)2); + //__sad((int)1, (int)2, 0); + //__uhadd((unsigned int)1, (unsigned int)3); + //__umul24((unsigned int)1, (unsigned int)2); + //__umul64hi((unsigned long long)1, (unsigned long long)2); + //__umulhi((unsigned int)1, (unsigned int)2); + //__urhadd((unsigned int)1, (unsigned int)2); + //__usad((unsigned int)1, (unsigned int)2, 0); } __global__ void compileIntegerIntrinsics(hipLaunchParm lp, int ignored) From ca142c6d9cf78570a3661918c60ba31549f89542 Mon Sep 17 00:00:00 2001 From: scchan Date: Mon, 1 Feb 2016 23:55:31 -0600 Subject: [PATCH 10/16] adding shfl, shfl_up, shfl_down, shfl_xor intrinsics [ROCm/hip commit: 04f3e3e59887a0098224513ed3011d689e1dd272] --- projects/hip/include/hcc_detail/hip_runtime.h | 41 +++++++++++++++++++ 1 file changed, 41 insertions(+) diff --git a/projects/hip/include/hcc_detail/hip_runtime.h b/projects/hip/include/hcc_detail/hip_runtime.h index 688c9e6959..734596a476 100644 --- a/projects/hip/include/hcc_detail/hip_runtime.h +++ b/projects/hip/include/hcc_detail/hip_runtime.h @@ -351,6 +351,47 @@ __device__ inline unsigned long long int __ballot( int input) return hc::__ballot( input); } +// warp shuffle functions +__device__ int __shfl(int input, int lane, int width) +{ + return hc::__shfl(input,lane,width); +} + +__device__ int __shfl_up(int input, unsigned int lane_delta, int width) +{ + return hc::__shfl_up(input,lane_delta,width); +} + +__device__ int __shfl_down(int input, unsigned int lane_delta, int width) +{ + return hc::__shfl_down(input,lane_delta,width); +} + +__device__ int __shfl_xor(int input, int lane_mask, int width) +{ + return hc::__shfl_xor(input,lane_mask,width); +} + +__device__ float __shfl(float input, int lane, int width) +{ + return hc::__shfl(input,lane,width); +} + +__device__ float __shfl_up(float input, unsigned int lane_delta, int width) +{ + return hc::__shfl_up(input,lane_delta,width); +} + +__device__ float __shfl_down(float input, unsigned int lane_delta, int width) +{ + return hc::__shfl_down(input,lane_delta,width); +} + +__device__ float __shfl_xor(float input, int lane_mask, int width) +{ + return hc::__shfl_xor(input,lane_mask,width); +} + #include // TODO: Choose whether default is precise math or fast math based on compilation flag. From c94a8de98317ddc25647dd6e787929160b2803a1 Mon Sep 17 00:00:00 2001 From: Jack Chung Date: Tue, 2 Feb 2016 16:27:42 +0800 Subject: [PATCH 11/16] Suppress linker warnings in case HCC distribution contains OpenCL/SPIR symbols [ROCm/hip commit: 1e7642c5699effe7e4699deac5550285fb981163] --- projects/hip/bin/hipcc | 2 ++ 1 file changed, 2 insertions(+) diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index ad9b8b162b..d001c6febe 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -66,6 +66,8 @@ if ($HIP_PLATFORM eq "hcc") { $HIPCXXFLAGS .= " -I$HIP_PATH/include/hcc_detail/cuda"; $HIPLDFLAGS = "-hc -L$HCC_HOME/lib -Wl,--rpath=$HCC_HOME/lib -lc++ -ldl -lpthread -Wl,--whole-archive -lmcwamp -Wl,--no-whole-archive"; + # Suppress linker warnings in case HCC distribution contains OpenCL/SPIR symbols + $HIPLDFLAGS .= " -Wl,--defsym=_binary_kernel_spir_end=0 -Wl,--defsym=_binary_kernel_spir_start=0 -Wl,--defsym=_binary_kernel_cl_start=0 -Wl,--defsym=_binary_kernel_cl_end=0"; $HIPLDFLAGS .= " -L$HSA_PATH/lib -lhsa-runtime64 -lhc_am"; # Add C++ libs for GCC. $HIPLDFLAGS .= " -lstdc++"; From 19d0b8095f86edc0f99dfd66ff0b1505917a19ab Mon Sep 17 00:00:00 2001 From: streamhsa Date: Tue, 2 Feb 2016 14:50:55 +0530 Subject: [PATCH 12/16] Added test for ballot and removing HIP_FUNCTION from hipSampleAtomicsTest.cpp -sandeep [ROCm/hip commit: 2b02fe082f30619f549a40f587b55014fdb7d572] --- projects/hip/tests/src/CMakeLists.txt | 2 + .../hip/tests/src/hipSimpleAtomicsTest.cpp | 4 +- projects/hip/tests/src/hip_anyall.cpp | 4 +- projects/hip/tests/src/hip_ballot.cpp | 48 +++++++++++++++++++ projects/hip/tests/src/hip_brev.cpp | 15 ------ projects/hip/tests/src/hip_clz.cpp | 18 ------- projects/hip/tests/src/hip_ffs.cpp | 16 ------- projects/hip/tests/src/hip_popc.cpp | 13 ----- 8 files changed, 54 insertions(+), 66 deletions(-) create mode 100644 projects/hip/tests/src/hip_ballot.cpp diff --git a/projects/hip/tests/src/CMakeLists.txt b/projects/hip/tests/src/CMakeLists.txt index 8da01db50a..7e41618cc2 100644 --- a/projects/hip/tests/src/CMakeLists.txt +++ b/projects/hip/tests/src/CMakeLists.txt @@ -103,6 +103,7 @@ endmacro() #install (TARGETS hipMemset DESTINATION bin) #install (TARGETS hipEventRecord DESTINATION bin) +make_hip_executable (hip_ballot hip_ballot.cpp) make_hip_executable (hip_anyall hip_anyall.cpp) make_hip_executable (hip_popc hip_popc.cpp) make_hip_executable (hip_clz hip_clz.cpp) @@ -120,6 +121,7 @@ make_hip_executable (hipMathFunctionsDevice hipMathFunctions.cpp hipSinglePrecis make_hip_executable (hipIntrinsics hipMathFunctions.cpp hipSinglePrecisionIntrinsics.cpp hipDoublePrecisionIntrinsics.cpp hipIntegerIntrinsics.cpp) target_link_libraries(hipMathFunctionsHost m) +make_test(hip_ballot " " ) make_test(hip_anyall " " ) make_test(hip_popc " " ) make_test(hip_brev " " ) diff --git a/projects/hip/tests/src/hipSimpleAtomicsTest.cpp b/projects/hip/tests/src/hipSimpleAtomicsTest.cpp index 16d3a7a746..f492643e41 100644 --- a/projects/hip/tests/src/hipSimpleAtomicsTest.cpp +++ b/projects/hip/tests/src/hipSimpleAtomicsTest.cpp @@ -190,7 +190,7 @@ int computeGold(int *gpuData, const int len) return true; } -__global__ void HIP_FUNCTION(testKernel,int *g_odata) +__global__ void testKernel(hipLaunchParm lp,int *g_odata) { // access thread id const unsigned int tid = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; @@ -236,7 +236,7 @@ __global__ void HIP_FUNCTION(testKernel,int *g_odata) // Atomic XOR atomicXor(&g_odata[10], tid); } -HIP_FUNCTION_END + int main(int argc, char **argv) { diff --git a/projects/hip/tests/src/hip_anyall.cpp b/projects/hip/tests/src/hip_anyall.cpp index 97bf664f6a..e126541766 100644 --- a/projects/hip/tests/src/hip_anyall.cpp +++ b/projects/hip/tests/src/hip_anyall.cpp @@ -32,8 +32,8 @@ __global__ void { int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; - device_any[hipThreadIdx_x>>6] = __any(tid >77); - device_all[hipThreadIdx_x>>6] = __all(tid >77); + device_any[hipThreadIdx_x>>6] = __any(tid -77); + device_all[hipThreadIdx_x>>6] = __all(tid -77); } diff --git a/projects/hip/tests/src/hip_ballot.cpp b/projects/hip/tests/src/hip_ballot.cpp new file mode 100644 index 0000000000..41087ca56c --- /dev/null +++ b/projects/hip/tests/src/hip_ballot.cpp @@ -0,0 +1,48 @@ +#include + +#include +#define HIP_ASSERT(x) (assert((x)==hipSuccess)) + +__global__ void + gpu_ballot(hipLaunchParm lp, unsigned int* device_ballot, int Num_Warps_per_Block) +{ + + int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + const unsigned int warp_num = hipThreadIdx_x >> 6; + atomicAdd(&device_ballot[warp_num+hipBlockIdx_x*Num_Warps_per_Block],__popcll(__ballot(tid - 245))); + +} + + +int main(int argc, char *argv[]) +{ + + unsigned int Num_Threads_per_Block = 512; + unsigned int Num_Blocks_per_Grid = 1; + unsigned int Num_Warps_per_Block = Num_Threads_per_Block/64; + unsigned int Num_Warps_per_Grid = (Num_Threads_per_Block*Num_Blocks_per_Grid)/64; + unsigned int* host_ballot = (unsigned int*)malloc(Num_Warps_per_Grid*sizeof(unsigned int)); + unsigned int* device_ballot; + HIP_ASSERT(hipMalloc((void**)&device_ballot, Num_Warps_per_Grid*sizeof(unsigned int))); + + for (int i=0; i Date: Tue, 2 Feb 2016 15:05:46 +0530 Subject: [PATCH 13/16] ADDED Support for __ffs() and __ffsll() having signed input -sandeep [ROCm/hip commit: 23904df99bad23d1d0ee5d99b83cd69baea1956b] --- projects/hip/include/hcc_detail/hip_runtime.h | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/projects/hip/include/hcc_detail/hip_runtime.h b/projects/hip/include/hcc_detail/hip_runtime.h index 734596a476..62386a14a2 100644 --- a/projects/hip/include/hcc_detail/hip_runtime.h +++ b/projects/hip/include/hcc_detail/hip_runtime.h @@ -325,6 +325,16 @@ __device__ inline unsigned int __ffsll(unsigned long long int input) return hc::__lastbit_u32_u64( input)+1; } +__device__ inline unsigned int __ffs(int input) +{ + return hc::__lastbit_u32_s32( input)+1; +} + +__device__ inline unsigned int __ffsll(long long int input) +{ + return hc::__lastbit_u32_s64( input)+1; +} + __device__ inline unsigned int __brev( unsigned int input) { return hc::__bitrev_b32( input); From af8cc35552c65fd32df800c04dfa156c1743b85b Mon Sep 17 00:00:00 2001 From: streamhsa Date: Tue, 2 Feb 2016 15:25:42 +0530 Subject: [PATCH 14/16] Adjusted the value of __any as per CUDA -sandeep [ROCm/hip commit: 974d4919023ca6224f5ada55dba68688bde60c7f] --- projects/hip/include/hcc_detail/hip_runtime.h | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/projects/hip/include/hcc_detail/hip_runtime.h b/projects/hip/include/hcc_detail/hip_runtime.h index 62386a14a2..f874a704bc 100644 --- a/projects/hip/include/hcc_detail/hip_runtime.h +++ b/projects/hip/include/hcc_detail/hip_runtime.h @@ -353,7 +353,8 @@ __device__ inline int __all( int input) __device__ inline int __any( int input) { - return hc::__any( input); + if( hc::__any( input)!=0) return 1; + else return 0; } __device__ inline unsigned long long int __ballot( int input) From 57d88913f8129ea6ce96716d45ddaa3f72417858 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Tue, 2 Feb 2016 10:02:48 -0600 Subject: [PATCH 15/16] Remove warning on ballot/any/all and pop/clz. Since these are supported in HIP no reason to emit warnings. [ROCm/hip commit: 3898f6af3cc206bb4081557022f0625318b84858] --- projects/hip/bin/hipify | 10 +++++----- projects/hip/tests/src/CMakeLists.txt | 13 +++++++++++++ projects/hip/tests/src/specialFunc.cu | 11 +++++++++++ 3 files changed, 29 insertions(+), 5 deletions(-) create mode 100644 projects/hip/tests/src/specialFunc.cu diff --git a/projects/hip/bin/hipify b/projects/hip/bin/hipify index 3eecc12abb..e5f17b69d0 100755 --- a/projects/hip/bin/hipify +++ b/projects/hip/bin/hipify @@ -624,12 +624,12 @@ sub warnUnsupportedSpecialFunctions "__ldg", # Cross-lane and warp-vote instructions: - "__all", - "__any", - "__ballot", + #"__all", + #"__any", + #"__ballot", - "__popc", - "__clz", + #"__popc", + #"__clz", "__shfl", "__shfl_up", diff --git a/projects/hip/tests/src/CMakeLists.txt b/projects/hip/tests/src/CMakeLists.txt index 3f2396b315..c1df4f376f 100644 --- a/projects/hip/tests/src/CMakeLists.txt +++ b/projects/hip/tests/src/CMakeLists.txt @@ -75,6 +75,7 @@ macro (make_hip_executable exe cpp) endif() endmacro() + macro (make_test exe ) string (REPLACE " " "" smush_args ${ARGN}) set (testname ${exe}${smush_args}.tst) @@ -87,6 +88,16 @@ macro (make_test exe ) endmacro() +macro (make_hipify_test sourceFile ) + #string (REPLACE " " "" smush_args ${ARGN}) + set (testname ${sourceFile}${smush_args}.tst) + + add_test (NAME ${testname} + COMMAND ${HIP_PATH}/bin/hipify ${PROJECT_SOURCE_DIR}/${sourceFile} ${ARGN} + ) +endmacro() + + macro (make_test_matches exe match_string) string (REPLACE " " "" smush_args ${ARGN}) set (testname ${exe}${smush_args}.tst) @@ -133,3 +144,5 @@ make_test(hipGridLaunch " " ) make_test(hipMemcpy " " ) make_test(hipHcc " " ) + +make_hipify_test(specialFunc.cu ) diff --git a/projects/hip/tests/src/specialFunc.cu b/projects/hip/tests/src/specialFunc.cu new file mode 100644 index 0000000000..c5c1931024 --- /dev/null +++ b/projects/hip/tests/src/specialFunc.cu @@ -0,0 +1,11 @@ +//Test to ensure hipify runs correctly. +// Hipify may report warnings for some missing/unsupported functions + +void __global__ +test_kernel(float *A) +{ + int tid = blockIdx.x * blockDim.x + threadIdx.x; + + float a = __ballot(tid < 16); + float b = __shfl(tid < 16); +} From 39fb16bc5f305777ec7a582cef6238a485de7fff Mon Sep 17 00:00:00 2001 From: scchan Date: Tue, 2 Feb 2016 12:53:17 -0600 Subject: [PATCH 16/16] add inline attribute to shfl functions [ROCm/hip commit: 265c42500f139334a4a1e8d142302350c67de06a] --- projects/hip/include/hcc_detail/hip_runtime.h | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/projects/hip/include/hcc_detail/hip_runtime.h b/projects/hip/include/hcc_detail/hip_runtime.h index f874a704bc..c8b9b0ebd7 100644 --- a/projects/hip/include/hcc_detail/hip_runtime.h +++ b/projects/hip/include/hcc_detail/hip_runtime.h @@ -363,42 +363,42 @@ __device__ inline unsigned long long int __ballot( int input) } // warp shuffle functions -__device__ int __shfl(int input, int lane, int width) +__device__ inline int __shfl(int input, int lane, int width) { return hc::__shfl(input,lane,width); } -__device__ int __shfl_up(int input, unsigned int lane_delta, int width) +__device__ inline int __shfl_up(int input, unsigned int lane_delta, int width) { return hc::__shfl_up(input,lane_delta,width); } -__device__ int __shfl_down(int input, unsigned int lane_delta, int width) +__device__ inline int __shfl_down(int input, unsigned int lane_delta, int width) { return hc::__shfl_down(input,lane_delta,width); } -__device__ int __shfl_xor(int input, int lane_mask, int width) +__device__ inline int __shfl_xor(int input, int lane_mask, int width) { return hc::__shfl_xor(input,lane_mask,width); } -__device__ float __shfl(float input, int lane, int width) +__device__ inline float __shfl(float input, int lane, int width) { return hc::__shfl(input,lane,width); } -__device__ float __shfl_up(float input, unsigned int lane_delta, int width) +__device__ inline float __shfl_up(float input, unsigned int lane_delta, int width) { return hc::__shfl_up(input,lane_delta,width); } -__device__ float __shfl_down(float input, unsigned int lane_delta, int width) +__device__ inline float __shfl_down(float input, unsigned int lane_delta, int width) { return hc::__shfl_down(input,lane_delta,width); } -__device__ float __shfl_xor(float input, int lane_mask, int width) +__device__ inline float __shfl_xor(float input, int lane_mask, int width) { return hc::__shfl_xor(input,lane_mask,width); }