From cd2b1f1baf8fa688591d399c2d54627a033a7e87 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Tue, 8 May 2018 15:36:01 -0400 Subject: [PATCH 1/2] Add __assert_fail, __device_trap and hipErrorAssert for clang [ROCm/clr commit: eeedfc556753946e5ca27bc90783167c7fdaf23f] --- .../hipamd/include/hip/hcc_detail/hip_runtime.h | 16 ++++++++++++++++ .../clr/hipamd/include/hip/hip_runtime_api.h | 2 ++ 2 files changed, 18 insertions(+) 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 92f06e9174..e9551d4f9a 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -612,6 +612,22 @@ extern const __device__ __attribute__((weak)) __hip_builtin_gridDim_t gridDim; #define hipGridDim_y gridDim.y #define hipGridDim_z gridDim.z +#pragma push_macro("__DEVICE__") +#define __DEVICE__ extern "C" __device__ inline __attribute__((always_inline)) \ + __attribute__((weak)) + +__DEVICE__ void __device_trap() __asm("llvm.trap"); + +__DEVICE__ void __assert_fail(const char * __assertion, + const char *__file, + unsigned int __line, + const char *__function) +{ + // Ignore all the args for now. + __device_trap(); +} +#pragma push_macro("__DEVICE__") + #endif #endif // HIP_HCC_DETAIL_RUNTIME_H diff --git a/projects/clr/hipamd/include/hip/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hip_runtime_api.h index bfde2e942d..2b36f3e140 100644 --- a/projects/clr/hipamd/include/hip/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/hip_runtime_api.h @@ -243,6 +243,8 @@ typedef enum __HIP_NODISCARD hipError_t { 1062, ///< Produced when trying to unlock a non-page-locked memory. hipErrorMapBufferObjectFailed = 1071, ///< Produced when the IPC memory attach failed from ROCr. + hipErrorAssert = + 1081, ///< Produced when the kernel calls assert. hipErrorTbd ///< Marker that more error codes are needed. } hipError_t; From 7d7125195caf823e7c75b0459b5db1f64fe8aa8f Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Tue, 8 May 2018 16:38:50 -0400 Subject: [PATCH 2/2] Fix warning about inlined function is not defined [ROCm/clr commit: 4c2d2b9e69cebf2a9742b7739d365ead3629177d] --- .../clr/hipamd/include/hip/hcc_detail/hip_runtime.h | 10 +++++----- 1 file changed, 5 insertions(+), 5 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 e9551d4f9a..e1b334aec3 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -613,15 +613,15 @@ extern const __device__ __attribute__((weak)) __hip_builtin_gridDim_t gridDim; #define hipGridDim_z gridDim.z #pragma push_macro("__DEVICE__") -#define __DEVICE__ extern "C" __device__ inline __attribute__((always_inline)) \ +#define __DEVICE__ extern "C" __device__ __attribute__((always_inline)) \ __attribute__((weak)) __DEVICE__ void __device_trap() __asm("llvm.trap"); -__DEVICE__ void __assert_fail(const char * __assertion, - const char *__file, - unsigned int __line, - const char *__function) +__DEVICE__ void inline __assert_fail(const char * __assertion, + const char *__file, + unsigned int __line, + const char *__function) { // Ignore all the args for now. __device_trap();