diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_assert.h b/projects/clr/hipamd/include/hip/amd_detail/hip_assert.h index 1ce6abb324..ba8e67bd91 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_assert.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_assert.h @@ -24,6 +24,14 @@ THE SOFTWARE. #if defined(__clang__) && defined(__HIP__) +// Use constexpr only in C++14 and later where constexpr void is allowed. +// In C++11, void is not a literal type and cannot be used as constexpr return type. +#if defined(__cplusplus) && __cplusplus >= 201402L + #define __HIP_CONSTEXPR_VOID constexpr +#else + #define __HIP_CONSTEXPR_VOID +#endif + // abort extern "C" __device__ inline __attribute__((weak)) void abort() { __builtin_trap(); } @@ -33,13 +41,13 @@ extern "C" __device__ inline __attribute__((weak)) void abort() { __builtin_trap // allows the function to exist as a global although its definition is // included in every compilation unit. #if defined(_WIN32) || defined(_WIN64) -extern "C" __device__ __attribute__((noinline)) __attribute__((weak)) void _wassert( +extern "C" __device__ __attribute__((noinline)) __attribute__((weak)) __HIP_CONSTEXPR_VOID void _wassert( const wchar_t* _msg, const wchar_t* _file, unsigned _line) { // FIXME: Need `wchar_t` support to generate assertion message. __builtin_trap(); } #else /* defined(_WIN32) || defined(_WIN64) */ -extern "C" __device__ __attribute__((noinline)) __attribute__((weak)) void __assert_fail( +extern "C" __device__ __attribute__((noinline)) __attribute__((weak)) __HIP_CONSTEXPR_VOID void __assert_fail( const char* assertion, const char* file, unsigned int line, const char* function) { const char fmt[] = "%s:%u: %s: Device-side assertion `%s' failed.\n"; @@ -77,7 +85,7 @@ extern "C" __device__ __attribute__((noinline)) __attribute__((weak)) void __ass __builtin_trap(); } -extern "C" __device__ __attribute__((noinline)) __attribute__((weak)) void __assertfail() { +extern "C" __device__ __attribute__((noinline)) __attribute__((weak)) __HIP_CONSTEXPR_VOID void __assertfail() { // ignore all the args for now. __builtin_trap(); } diff --git a/projects/hip-tests/catch/unit/assertion/CMakeLists.txt b/projects/hip-tests/catch/unit/assertion/CMakeLists.txt index 6e2ef391ce..77d8a2401a 100644 --- a/projects/hip-tests/catch/unit/assertion/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/assertion/CMakeLists.txt @@ -30,6 +30,7 @@ elseif(HIP_PLATFORM MATCHES "amd") set(TEST_SRC static_assert.cc assert.cc + constexpr_assert.cc ) hip_add_exe_to_target(NAME AssertionTest TEST_SRC ${TEST_SRC} diff --git a/projects/hip-tests/catch/unit/assertion/constexpr_assert.cc b/projects/hip-tests/catch/unit/assertion/constexpr_assert.cc new file mode 100644 index 0000000000..7671ef3dd8 --- /dev/null +++ b/projects/hip-tests/catch/unit/assertion/constexpr_assert.cc @@ -0,0 +1,59 @@ +/* +Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include + +/** + * @addtogroup constexpr_assert constexpr_assert + * @{ + * @ingroup DeviceLanguageTest + * Tests for using assert() inside constexpr device functions. + */ + +__device__ constexpr int safe_divide(int a, int b) { + assert(b != 0); + return a / b; +} + +__device__ constexpr int result = safe_divide(10, 2); + +/** + * Test Description + * ------------------------ + * - Verifies that assert() can be used in constexpr __device__ functions. + * - Tests compile-time evaluation of constexpr device function with assert. + * - The assert condition is true, so it should not trigger. + * Test source + * ------------------------ + * - unit/assertion/constexpr_assert.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.3 + */ +TEST_CASE("Unit_ConstexprAssert_Positive_SafeDivide") { + static_assert(result == 5, "safe_divide(10, 2) should equal 5"); + REQUIRE(result == 5); +} + +/** + * End doxygen group DeviceLanguageTest. + * @} + */ +