From c9f64bbd2d06bbc7c348aee66bdebd09620bebc7 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Fri, 17 Mar 2017 10:25:56 -0500 Subject: [PATCH] Add simple device-side assert macro Currently swallows asserts but will compile. --- hipamd/include/hip/hcc_detail/hip_runtime.h | 11 +++++++++++ hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp | 2 ++ 2 files changed, 13 insertions(+) diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h index 332e9bab46..673463635f 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -77,6 +77,17 @@ extern int HIP_TRACE_API; #define __HCC_ACCELERATOR__ __KALMAR_ACCELERATOR__ #endif + + + +// TODO-HCC add a dummy implementation of assert, need to replace with a proper kernel exit call. +#if __HIP_DEVICE_COMPILE__ == 1 + #undef assert + #define assert(COND) { if (COND) {} } +#endif + + + // Feature tests: #if defined(__HCC_ACCELERATOR__) && (__HCC_ACCELERATOR__ != 0) // Device compile and not host compile: diff --git a/hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp b/hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp index 6bf13a0809..27af03ced3 100644 --- a/hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp +++ b/hipamd/tests/src/deviceLib/hipIntegerIntrinsics.cpp @@ -59,6 +59,8 @@ __device__ void integer_intrinsics() __umulhi((unsigned int)1, (unsigned int)2); __urhadd((unsigned int)1, (unsigned int)2); __usad((unsigned int)1, (unsigned int)2, 0); + + assert(1); } __global__ void compileIntegerIntrinsics(hipLaunchParm lp, int ignored)