From b0602d501a199f4c9c6e04e78f884d693c46fdb7 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/hip commit: cf12a9c049cf9254cd7d1f0e62cee161d71820d7] --- projects/hip/include/hip/hcc_detail/hip_fp16.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_fp16.h b/projects/hip/include/hip/hcc_detail/hip_fp16.h index 68f0e35f5f..5a6e650069 100644 --- a/projects/hip/include/hip/hcc_detail/hip_fp16.h +++ b/projects/hip/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 6b8913d4e59bdf30ce82e7a582091fa5eb267df0 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/hip commit: 0b211bb4292954230d9576ac50b9491e1ef7def7] --- projects/hip/tests/src/deviceLib/hipTestIncludeMath.cpp | 9 +-------- 1 file changed, 1 insertion(+), 8 deletions(-) diff --git a/projects/hip/tests/src/deviceLib/hipTestIncludeMath.cpp b/projects/hip/tests/src/deviceLib/hipTestIncludeMath.cpp index 6063eee76c..92b93eda36 100644 --- a/projects/hip/tests/src/deviceLib/hipTestIncludeMath.cpp +++ b/projects/hip/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 fc3c4aca009cf96d77f59ed0ced4290379261502 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/hip commit: cef5261fa9e1e63672d2857e22c63839b629aea0] --- projects/hip/include/hip/nvcc_detail/hip_runtime.h | 3 +++ 1 file changed, 3 insertions(+) diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime.h b/projects/hip/include/hip/nvcc_detail/hip_runtime.h index fe0b19bf8c..19d740a1ee 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime.h @@ -118,6 +118,9 @@ typedef int hipLaunchParm; } #endif +#define __clock() clock() +#define __clock64() clock64() + #endif #endif From e14fc1c65e8693fcfaeac7a020760748492e897e 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/hip commit: 94a0589eb5de05a05c6ae8e3295b7a4b418face9] --- projects/hip/tests/src/deviceLib/hipSimpleAtomicsTest.cpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/projects/hip/tests/src/deviceLib/hipSimpleAtomicsTest.cpp b/projects/hip/tests/src/deviceLib/hipSimpleAtomicsTest.cpp index 129d7c1aa8..ec38a5d327 100644 --- a/projects/hip/tests/src/deviceLib/hipSimpleAtomicsTest.cpp +++ b/projects/hip/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 5947c0e94dfb382c8a1d4744f901c78446d7ac68 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/hip commit: cca2c5afc80818601a338b25c2c353a31f65a05e] --- .../hip/tests/src/deviceLib/hipTestClock.cpp | 5 +---- .../hip/tests/src/deviceLib/hipTestNew.cpp | 4 +--- .../tests/src/kernel/hipTestGlobalVariable.cpp | 18 ++++++------------ 3 files changed, 8 insertions(+), 19 deletions(-) diff --git a/projects/hip/tests/src/deviceLib/hipTestClock.cpp b/projects/hip/tests/src/deviceLib/hipTestClock.cpp index 46f64e35a3..ee6dca8a42 100644 --- a/projects/hip/tests/src/deviceLib/hipTestClock.cpp +++ b/projects/hip/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/hip/tests/src/deviceLib/hipTestNew.cpp b/projects/hip/tests/src/deviceLib/hipTestNew.cpp index 60774ff21d..d644f8b483 100644 --- a/projects/hip/tests/src/deviceLib/hipTestNew.cpp +++ b/projects/hip/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/hip/tests/src/kernel/hipTestGlobalVariable.cpp b/projects/hip/tests/src/kernel/hipTestGlobalVariable.cpp index 2209e2c254..8ab8bef9c2 100644 --- a/projects/hip/tests/src/kernel/hipTestGlobalVariable.cpp +++ b/projects/hip/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 1aca8f7d9d0375815420aef468d5578ebd368b39 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/hip commit: 30acc99e134314e43e77b72fa6dc10a5fdc3acb9] --- projects/hip/install.sh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/install.sh b/projects/hip/install.sh index f8ad640798..c7dc48e1a8 100755 --- a/projects/hip/install.sh +++ b/projects/hip/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 694b6e105a02ecae0822cc4271ecc69eccd98c62 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/hip commit: fc228c7ea6f36b8a3ec75505ae1500be02606635] --- projects/hip/include/hip/hcc_detail/hip_runtime.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime.h b/projects/hip/include/hip/hcc_detail/hip_runtime.h index e5f0fb52fa..d8093e6646 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime.h +++ b/projects/hip/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 3d60af4f35426b6f354d4cc6ff74f5e97d9d9704 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/hip commit: cdfd82f1dee9b180b46bd256b81184600532c53a] --- projects/hip/include/hip/hcc_detail/hip_memory.h | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/projects/hip/include/hip/hcc_detail/hip_memory.h b/projects/hip/include/hip/hcc_detail/hip_memory.h index 2394a05d0f..739e488ca3 100644 --- a/projects/hip/include/hip/hcc_detail/hip_memory.h +++ b/projects/hip/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