From 1d0e3c3ece6a2d46547ddb217b9b91065f53ed10 Mon Sep 17 00:00:00 2001 From: Johannes M Dieterich Date: Thu, 13 Sep 2018 13:55:43 -0500 Subject: [PATCH 1/8] Only LLVM6 and higher contain the necessary intrinsics. [ROCm/clr commit: dbe0cbcb93d8120ae1f2d85a2e9bc3ab999f8e0c] --- projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h index 68f0e35f5f..5a6e650069 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h @@ -29,7 +29,7 @@ THE SOFTWARE. #include #endif -#if defined(__clang__) && (__clang_major__ > 3) +#if defined(__clang__) && (__clang_major__ > 5) typedef _Float16 _Float16_2 __attribute__((ext_vector_type(2))); struct __half_raw { From 51f5e9ea04b3fd8995429c46bfa26c838295af25 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 17 Sep 2018 15:15:49 +0530 Subject: [PATCH 2/8] [dtests] Fix deviceLib/hipTestIncludeMath on nvcc path Change-Id: Ifb121886bec6c1134e98bd194ff1713f28454d33 [ROCm/clr commit: 644fc85a37ae2ed0b61516d935b44da9afa0e4ed] --- .../hipamd/tests/src/deviceLib/hipTestIncludeMath.cpp | 9 +-------- 1 file changed, 1 insertion(+), 8 deletions(-) diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipTestIncludeMath.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipTestIncludeMath.cpp index 6063eee76c..92b93eda36 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipTestIncludeMath.cpp +++ b/projects/clr/hipamd/tests/src/deviceLib/hipTestIncludeMath.cpp @@ -30,6 +30,7 @@ THE SOFTWARE. // Incorrect implementation causes compilation failure due to conflict // declartions. +#include #include // Test __HIP_DEVICE_COMPILE__ is defined after math_functions.h @@ -45,14 +46,6 @@ __device__ __host__ inline void throw_std_bad_alloc() #endif } -// Test __HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__ and __HIP_ARCH_HAS_DYNAMIC_PARALLEL__ -// is defined. Eigen HIP/hcc/Half.h __ldg depends on this. -#if !defined(__HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__) || \ - !defined(__HIP_ARCH_HAS_DYNAMIC_PARALLEL__) -#error \ - "__HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__ or __HIP_ARCH_HAS_DYNAMIC_PARALLEL__ not defined" -#endif - #include #include "test_common.h" From 9623689bc736c286f5ce24e0751823687e8d36c6 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 17 Sep 2018 15:23:30 +0530 Subject: [PATCH 3/8] Add mappings for __clock* in nvcc_detail/hip_runtime.h Change-Id: Ibcecf52f3e69298268d921efc036090544fa0ed0 [ROCm/clr commit: 933c13f00308c368de576f7080831db23e6f425c] --- projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime.h | 3 +++ 1 file changed, 3 insertions(+) diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime.h index fe0b19bf8c..19d740a1ee 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime.h @@ -118,6 +118,9 @@ typedef int hipLaunchParm; } #endif +#define __clock() clock() +#define __clock64() clock64() + #endif #endif From 6e7f871105f21fc148e8d863f87860f0b61e8720 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 17 Sep 2018 15:26:45 +0530 Subject: [PATCH 4/8] [dtests] Fix deviceLib/hipSimpleAtomicsTest for nvcc Change-Id: I23cbea2820d41da6f6a1bcab4bc3b59ac21799dd [ROCm/clr commit: 706614776021e05b936dd99c761ab31a590f38b7] --- .../clr/hipamd/tests/src/deviceLib/hipSimpleAtomicsTest.cpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipSimpleAtomicsTest.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipSimpleAtomicsTest.cpp index 129d7c1aa8..ec38a5d327 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipSimpleAtomicsTest.cpp +++ b/projects/clr/hipamd/tests/src/deviceLib/hipSimpleAtomicsTest.cpp @@ -18,7 +18,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../test_common.cpp + * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -std=c++11 --gpu-architecture=sm_60 * RUN: %t * HIT_END */ @@ -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); @@ -333,4 +334,4 @@ int main(int argc, char** argv) { hipDeviceReset(); printf("%s completed, returned %s\n", sampleName, testResult ? "OK" : "ERROR!"); exit(testResult ? EXIT_SUCCESS : EXIT_FAILURE); -} \ No newline at end of file +} From 0bae7eaad33a5ccadbbbb11e85167362d648f955 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 17 Sep 2018 15:32:05 +0530 Subject: [PATCH 5/8] [dests] Fix hipTestClock, hipTestNew & hipTestGlobalVariable tests for nvcc nvcc does not support global kernels in struct/class Change-Id: I2d7297e0c3725564215e20dbdd31c0bb8d7a07de [ROCm/clr commit: f63ffaf6e5c91dcbe9301cc4ae6e32f33c30eaf5] --- .../tests/src/deviceLib/hipTestClock.cpp | 5 +---- .../hipamd/tests/src/deviceLib/hipTestNew.cpp | 4 +--- .../tests/src/kernel/hipTestGlobalVariable.cpp | 18 ++++++------------ 3 files changed, 8 insertions(+), 19 deletions(-) diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipTestClock.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipTestClock.cpp index 46f64e35a3..ee6dca8a42 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipTestClock.cpp +++ b/projects/clr/hipamd/tests/src/deviceLib/hipTestClock.cpp @@ -33,8 +33,6 @@ THE SOFTWARE. #define LEN 512 #define SIZE 2048 -struct TestClock { - static __global__ void kernel1(int* Ad) { int tid = threadIdx.x + blockIdx.x * blockDim.x; Ad[tid] = clock() + clock64() + __clock() + __clock64(); @@ -61,9 +59,8 @@ struct TestClock { assert(0 != A[i]); } } -}; int main() { - TestClock().run(); + run(); passed(); } diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipTestNew.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipTestNew.cpp index 60774ff21d..d644f8b483 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipTestNew.cpp +++ b/projects/clr/hipamd/tests/src/deviceLib/hipTestNew.cpp @@ -33,7 +33,6 @@ THE SOFTWARE. #define LEN 512 #define SIZE 2048 -struct TestPlacementNew { class A { public: __device__ A() { @@ -63,9 +62,8 @@ struct TestPlacementNew { assert(i == A[i]); } } -}; int main() { - TestPlacementNew().run(); + run(); passed(); } diff --git a/projects/clr/hipamd/tests/src/kernel/hipTestGlobalVariable.cpp b/projects/clr/hipamd/tests/src/kernel/hipTestGlobalVariable.cpp index 2209e2c254..8ab8bef9c2 100644 --- a/projects/clr/hipamd/tests/src/kernel/hipTestGlobalVariable.cpp +++ b/projects/clr/hipamd/tests/src/kernel/hipTestGlobalVariable.cpp @@ -33,15 +33,14 @@ THE SOFTWARE. #define LEN 512 #define SIZE 2048 -struct TestConstantGlobalVar { - static __constant__ int ConstantGlobalVar; + __constant__ int ConstantGlobalVar = 123; static __global__ void kernel(int* Ad) { int tid = threadIdx.x + blockIdx.x * blockDim.x; Ad[tid] = ConstantGlobalVar; } - void run() { + void runTestConstantGlobalVar() { int *A, *Ad; A = new int[LEN]; for (unsigned i = 0; i < LEN; i++) { @@ -56,11 +55,8 @@ struct TestConstantGlobalVar { assert(123 == A[i]); } } -}; -__constant__ int TestConstantGlobalVar::ConstantGlobalVar = 123; -struct TestGlobalArray { - static __device__ int GlobalArray[LEN]; + __device__ int GlobalArray[LEN]; static __global__ void kernelWrite() { int tid = threadIdx.x + blockIdx.x * blockDim.x; @@ -71,7 +67,7 @@ struct TestGlobalArray { Ad[tid] = GlobalArray[tid]; } - void run() { + void runTestGlobalArray() { int *A, *Ad; A = new int[LEN]; for (unsigned i = 0; i < LEN; i++) { @@ -87,11 +83,9 @@ struct TestGlobalArray { assert(i == A[i]); } } -}; -__device__ int TestGlobalArray::GlobalArray[LEN]; int main() { - TestConstantGlobalVar().run(); - TestGlobalArray().run(); + runTestConstantGlobalVar(); + runTestGlobalArray(); passed(); } From 564a45497b29ce6812a55eb68bfc4f5a42107373 Mon Sep 17 00:00:00 2001 From: VincentSC Date: Mon, 17 Sep 2018 13:25:12 +0200 Subject: [PATCH 6/8] Update install.sh 'rename' added as a Linux prerequisite. [ROCm/clr commit: fcc9aa3412efacdec29f0ecedaa24d2bda6c4c37] --- projects/clr/hipamd/install.sh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/install.sh b/projects/clr/hipamd/install.sh index f8ad640798..c7dc48e1a8 100755 --- a/projects/clr/hipamd/install.sh +++ b/projects/clr/hipamd/install.sh @@ -25,7 +25,7 @@ popd () { function setupENV() { sudo apt-get update - sudo apt-get install dpkg-dev rpm doxygen libelf-dev + sudo apt-get install dpkg-dev rpm doxygen libelf-dev rename } function buildHIP() From eb9c852dc57f6a41059eb2013b8cebb5a59db3d6 Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Mon, 17 Sep 2018 11:19:35 -0400 Subject: [PATCH 7/8] Fix hipLaunchKernelGGL for hip-clang Do not decay function pointer type of the kernel argument passed to hipLaunchKernelGGL and hipLaunchKernel, otherwise some type information is lost which may cause type inference failure for the template. This issue caused compilation error of FeatureLPPooling in Caffe2/PyTorch and this patch fixes that. [ROCm/clr commit: 0471ce2da65b48fe4504906f556ec86bc3ce0d63] --- projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h | 4 ++-- 1 file changed, 2 insertions(+), 2 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 e5f0fb52fa..d8093e6646 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -334,13 +334,13 @@ extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, gri typedef int hipLaunchParm; template -inline void hipLaunchKernelGGL(F kernelName, const dim3& numblocks, const dim3& numthreads, +inline void hipLaunchKernelGGL(F&& kernelName, const dim3& numblocks, const dim3& numthreads, unsigned memperblock, hipStream_t streamId, Args... args) { kernelName<<>>(args...); } template -inline void hipLaunchKernel(F kernel, const dim3& numBlocks, const dim3& dimBlocks, +inline void hipLaunchKernel(F&& kernel, const dim3& numBlocks, const dim3& dimBlocks, std::uint32_t groupMemBytes, hipStream_t stream, Args... args) { hipLaunchKernelGGL(kernel, numBlocks, dimBlocks, groupMemBytes, stream, hipLaunchParm{}, std::move(args)...); From c0592f16e602f6c213122acfca28ffd0c90113dd Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Mon, 17 Sep 2018 16:50:42 -0400 Subject: [PATCH 8/8] Disable device code for gcc in hip_memory.h These device code should only be seen by HCC or hip-clang. They causd build failure for HIP-VDI runtime and should be disabled for gcc. [ROCm/clr commit: b23db4e45244bcb8c8085fe6347d2aa9b6c2ccb8] --- projects/clr/hipamd/include/hip/hcc_detail/hip_memory.h | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_memory.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_memory.h index 2394a05d0f..739e488ca3 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_memory.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_memory.h @@ -39,6 +39,8 @@ THE SOFTWARE. #define __HIP_SIZE_OF_HEAP (__HIP_NUM_PAGES * __HIP_SIZE_OF_PAGE) +#if __HCC__ || __HIP__ + __attribute__((weak)) __device__ char __hip_device_heap[__HIP_SIZE_OF_HEAP]; __attribute__((weak)) __device__ uint32_t __hip_device_page_flag[__HIP_NUM_PAGES]; @@ -102,4 +104,6 @@ extern "C" inline __device__ void* __hip_free(void* ptr) { return nullptr; } +#endif + #endif // HIP_INCLUDE_HIP_HCC_DETAIL_HIP_MEMORY_H