From 96f6b6e2512bddcdd783ab62be930cb1712a372d Mon Sep 17 00:00:00 2001 From: Shadi Dashmiz <94885391+shadidashmiz@users.noreply.github.com> Date: Wed, 17 Dec 2025 11:15:20 -0500 Subject: [PATCH] SWDEV-571304 : Fix the constructor for __half (#2240) - comply with cuda - Fix usecase for constexpr Signed-off-by: sdashmiz --- .../include/hip/amd_detail/amd_hip_fp16.h | 2 +- .../catch/unit/deviceLib/CMakeLists.txt | 1 + .../unit/deviceLib/hipTestHalfConstexpr.cc | 126 ++++++++++++++++++ 3 files changed, 128 insertions(+), 1 deletion(-) create mode 100644 projects/hip-tests/catch/unit/deviceLib/hipTestHalfConstexpr.cc diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h index 21c69640cb..7e67cee654 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_fp16.h @@ -99,7 +99,7 @@ struct __half { // CREATORS __HOST_DEVICE__ __half() = default; - __HOST_DEVICE__ constexpr __half(const __half_raw& x) : data{x.data} {} + __HOST_DEVICE__ constexpr __half(const __half_raw& x) : __x{x.x} {} #if !defined(__HIP_NO_HALF_CONVERSIONS__) __HOST_DEVICE__ __half(decltype(data) x) : data{x} {} diff --git a/projects/hip-tests/catch/unit/deviceLib/CMakeLists.txt b/projects/hip-tests/catch/unit/deviceLib/CMakeLists.txt index fe079b22ee..bbe5597b09 100644 --- a/projects/hip-tests/catch/unit/deviceLib/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/deviceLib/CMakeLists.txt @@ -77,6 +77,7 @@ set(AMD_TEST_SRC hipVectorTypesHost.cc hipVectorTypesDevice.cc hipTestHalf.cc + hipTestHalfConstexpr.cc hipComplex.cc hipTestFMA.cc hipTestNativeHalf.cc diff --git a/projects/hip-tests/catch/unit/deviceLib/hipTestHalfConstexpr.cc b/projects/hip-tests/catch/unit/deviceLib/hipTestHalfConstexpr.cc new file mode 100644 index 0000000000..f1331fbe9d --- /dev/null +++ b/projects/hip-tests/catch/unit/deviceLib/hipTestHalfConstexpr.cc @@ -0,0 +1,126 @@ +/* +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 +#include + + +constexpr unsigned short kHalfTwo = 0x4000U; // 2.0 +constexpr unsigned short kHalfPi = 0x4248U; // ~3.140625 (closest to pi) +constexpr unsigned short kHalfInf = 0x7C00U; // +Infinity + +/** + * Constexpr function to create __half from bit pattern. + * This pattern is used by libhipcxx for mathematical constants. + */ +__host__ __device__ constexpr __half makeHalfFromBits(unsigned short bits) { + return __half{__half_raw{.x = bits}}; +} + +// Constexpr half-precision constants using the pattern from the bug report + +constexpr __half kConstTwo = makeHalfFromBits(kHalfTwo); +constexpr __half kConstPi = makeHalfFromBits(kHalfPi); +constexpr __half kConstInf = makeHalfFromBits(kHalfInf); + + +/** + * Device kernel that uses constexpr __half values. + * Tests that constexpr __half can be used in device code. + */ +__global__ void testConstexprHalfDevice(float* results) { + + // Verify the constexpr values are usable in device code + results[0] = __half2float(kConstTwo); + results[1]= __half2float(kConstPi); + + // Test arithmetic with constexpr values + results[2] = __half2float(__hmul(kConstTwo, kConstPi)); // 2 * pi +} + +/** + * Test Description + * ------------------------ + * - Tests constexpr construction of __half from __half_raw{.x = bits} + * - This is the pattern used by libhipcxx for mathematical constants + * - Verifies fix for union active member issue in constexpr evaluation + */ +TEST_CASE("Unit_hipTestHalfConstexpr_HostConstexpr") { + // Test that constexpr __half values have correct bit patterns on host + + SECTION("Two") { + constexpr __half two = makeHalfFromBits(kHalfTwo); + float f = __half2float(two); + REQUIRE(f == 2.0f); + REQUIRE(__half_as_ushort(two) == kHalfTwo); + } + + SECTION("Pi approximation") { + constexpr __half pi = makeHalfFromBits(kHalfPi); + float f = __half2float(pi); + // Half precision pi is approximately 3.140625 + REQUIRE(f == Catch::Approx(3.14159f).epsilon(0.01)); + REQUIRE(__half_as_ushort(pi) == kHalfPi); + } + + SECTION("Infinity") { + constexpr __half inf = makeHalfFromBits(kHalfInf); + REQUIRE(__hisinf(inf)); + REQUIRE(__half_as_ushort(inf) == kHalfInf); + } + + SECTION("File scope constexpr values") { + // Test that file-scope constexpr values are correct + REQUIRE(__half2float(kConstTwo) == 2.0f); + } +} + +/** + * Test Description + * ------------------------ + * - Tests that constexpr __half values can be used in device kernels + * - Verifies both file-scope and kernel-local constexpr values work + */ +TEST_CASE("Unit_hipTestHalfConstexpr_DeviceConstexpr") { + constexpr size_t numResults = 3; + float* results_d = nullptr; + std::vector results_h(numResults, 0.0f); + + HIP_CHECK(hipMalloc(&results_d, numResults * sizeof(float))); + + testConstexprHalfDevice<<<1, 1>>>(results_d); + HIP_CHECK(hipDeviceSynchronize()); + + HIP_CHECK(hipMemcpy(results_h.data(), results_d, numResults * sizeof(float), + hipMemcpyDeviceToHost)); + + // Verify results + REQUIRE(results_h[0] == 2.0f); // kConstTwo + REQUIRE(results_h[1] == Catch::Approx(3.14159f).epsilon(0.01)); // kConstPi + REQUIRE(results_h[2] == Catch::Approx(6.28f).epsilon(0.01)); // 2 * pi + + HIP_CHECK(hipFree(results_d)); +} + +