From 0b211bb4292954230d9576ac50b9491e1ef7def7 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 17 Sep 2018 15:15:49 +0530 Subject: [PATCH 1/4] [dtests] Fix deviceLib/hipTestIncludeMath on nvcc path Change-Id: Ifb121886bec6c1134e98bd194ff1713f28454d33 --- tests/src/deviceLib/hipTestIncludeMath.cpp | 9 +-------- 1 file changed, 1 insertion(+), 8 deletions(-) diff --git a/tests/src/deviceLib/hipTestIncludeMath.cpp b/tests/src/deviceLib/hipTestIncludeMath.cpp index 6063eee76c..92b93eda36 100644 --- a/tests/src/deviceLib/hipTestIncludeMath.cpp +++ b/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 cef5261fa9e1e63672d2857e22c63839b629aea0 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 17 Sep 2018 15:23:30 +0530 Subject: [PATCH 2/4] Add mappings for __clock* in nvcc_detail/hip_runtime.h Change-Id: Ibcecf52f3e69298268d921efc036090544fa0ed0 --- include/hip/nvcc_detail/hip_runtime.h | 3 +++ 1 file changed, 3 insertions(+) diff --git a/include/hip/nvcc_detail/hip_runtime.h b/include/hip/nvcc_detail/hip_runtime.h index fe0b19bf8c..19d740a1ee 100644 --- a/include/hip/nvcc_detail/hip_runtime.h +++ b/include/hip/nvcc_detail/hip_runtime.h @@ -118,6 +118,9 @@ typedef int hipLaunchParm; } #endif +#define __clock() clock() +#define __clock64() clock64() + #endif #endif From 94a0589eb5de05a05c6ae8e3295b7a4b418face9 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 17 Sep 2018 15:26:45 +0530 Subject: [PATCH 3/4] [dtests] Fix deviceLib/hipSimpleAtomicsTest for nvcc Change-Id: I23cbea2820d41da6f6a1bcab4bc3b59ac21799dd --- tests/src/deviceLib/hipSimpleAtomicsTest.cpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/tests/src/deviceLib/hipSimpleAtomicsTest.cpp b/tests/src/deviceLib/hipSimpleAtomicsTest.cpp index 129d7c1aa8..ec38a5d327 100644 --- a/tests/src/deviceLib/hipSimpleAtomicsTest.cpp +++ b/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 cca2c5afc80818601a338b25c2c353a31f65a05e Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 17 Sep 2018 15:32:05 +0530 Subject: [PATCH 4/4] [dests] Fix hipTestClock, hipTestNew & hipTestGlobalVariable tests for nvcc nvcc does not support global kernels in struct/class Change-Id: I2d7297e0c3725564215e20dbdd31c0bb8d7a07de --- tests/src/deviceLib/hipTestClock.cpp | 5 +---- tests/src/deviceLib/hipTestNew.cpp | 4 +--- tests/src/kernel/hipTestGlobalVariable.cpp | 18 ++++++------------ 3 files changed, 8 insertions(+), 19 deletions(-) diff --git a/tests/src/deviceLib/hipTestClock.cpp b/tests/src/deviceLib/hipTestClock.cpp index 46f64e35a3..ee6dca8a42 100644 --- a/tests/src/deviceLib/hipTestClock.cpp +++ b/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/tests/src/deviceLib/hipTestNew.cpp b/tests/src/deviceLib/hipTestNew.cpp index 60774ff21d..d644f8b483 100644 --- a/tests/src/deviceLib/hipTestNew.cpp +++ b/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/tests/src/kernel/hipTestGlobalVariable.cpp b/tests/src/kernel/hipTestGlobalVariable.cpp index 2209e2c254..8ab8bef9c2 100644 --- a/tests/src/kernel/hipTestGlobalVariable.cpp +++ b/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(); }