From d1d115fdbbc44a483c6ff766a2d13f94e4caf668 Mon Sep 17 00:00:00 2001 From: Steven Chung Date: Fri, 25 Sep 2026 23:56:37 -0400 Subject: [PATCH 1/5] fix(miopen): resolve uint64_t typedef conflict with system headers on LP64 Linux miopen_cstdint.hpp typedefs uint64_t as __hip_internal::uint64_t (unsigned long long), but on LP64 Linux defines uint64_t as unsigned long. These are distinct C++ types despite being the same width, causing a hard compilation error when both definitions are visible during MIOpen's runtime kernel compilation (comgr/HIPRTC JIT path). Use platform-appropriate types: unsigned long on LP64 (Linux/macOS) and unsigned long long on LLP64 (Windows), matching what the system provides on each platform. Fixes: https://github.com/ROCm/TheRock/issues/6258 Co-Authored-By: Claude --- .../miopen/src/kernels/miopen_cstdint.hpp | 9 ++- .../miopen/test/gtest/miopen_cstdint_lp64.cpp | 68 +++++++++++++++++++ 2 files changed, 75 insertions(+), 2 deletions(-) create mode 100644 projects/miopen/test/gtest/miopen_cstdint_lp64.cpp diff --git a/projects/miopen/src/kernels/miopen_cstdint.hpp b/projects/miopen/src/kernels/miopen_cstdint.hpp index fa8d68f5fa58..cb1eff85bc31 100644 --- a/projects/miopen/src/kernels/miopen_cstdint.hpp +++ b/projects/miopen/src/kernels/miopen_cstdint.hpp @@ -33,8 +33,13 @@ typedef unsigned short uint16_t; #if HIP_PACKAGE_VERSION_FLAT >= 6000025000ULL typedef signed int int32_t; typedef unsigned int uint32_t; -typedef __hip_internal::uint64_t uint64_t; -typedef __hip_internal::int64_t int64_t; +#ifdef __LP64__ +typedef unsigned long uint64_t; +typedef signed long int64_t; +#else +typedef unsigned long long uint64_t; +typedef signed long long int64_t; +#endif #endif #else diff --git a/projects/miopen/test/gtest/miopen_cstdint_lp64.cpp b/projects/miopen/test/gtest/miopen_cstdint_lp64.cpp new file mode 100644 index 000000000000..2b8a2cf13926 --- /dev/null +++ b/projects/miopen/test/gtest/miopen_cstdint_lp64.cpp @@ -0,0 +1,68 @@ +/******************************************************************************* + * + * MIT License + * + * Copyright (c) 2026 Advanced Micro Devices, Inc. + * + * 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. + * + *******************************************************************************/ + +// Regression test for TheRock#6258 / LP64 typedef conflict in miopen_cstdint.hpp. +// +// On LP64 platforms (64-bit Linux, macOS), defines uint64_t as +// 'unsigned long', but __hip_internal::uint64_t is 'unsigned long long'. +// These are distinct C++ types despite being the same width, so a typedef +// redefinition is a hard error when both are visible. + +#include +#include +#include + +#define MIOPEN_HIP_RUNTIME_COMPILE +#define HIP_PACKAGE_VERSION_FLAT 7015026333ULL + +// Provide __hip_internal types as they would appear during HIPRTC compilation +namespace __hip_internal { +typedef unsigned long long uint64_t; +typedef signed long long int64_t; +} // namespace __hip_internal + +// Include the header under test — this must not conflict with above +#include "miopen_cstdint.hpp" + +#undef MIOPEN_HIP_RUNTIME_COMPILE + +TEST(MiopenCstdintLP64, Uint64TypeMatchesSystem) +{ + EXPECT_TRUE((std::is_same<::uint64_t, std::uint64_t>::value)) + << "miopen_cstdint.hpp uint64_t must be the same type as uint64_t"; +} + +TEST(MiopenCstdintLP64, Int64TypeMatchesSystem) +{ + EXPECT_TRUE((std::is_same<::int64_t, std::int64_t>::value)) + << "miopen_cstdint.hpp int64_t must be the same type as int64_t"; +} + +TEST(MiopenCstdintLP64, SizesAre64Bit) +{ + EXPECT_EQ(sizeof(::uint64_t), 8u); + EXPECT_EQ(sizeof(::int64_t), 8u); +} From 54ac41642ba69798a493fbff3b9f82e9a317fd97 Mon Sep 17 00:00:00 2001 From: Steven Chung Date: Sat, 26 Sep 2026 09:40:27 -0400 Subject: [PATCH 2/5] fix(miopen): use MIOpen gtest naming convention for test suite Rename test suites to CPU_MiopenCstdintLP64_NONE to match the required pattern: (CPU|GPU)__. Co-Authored-By: Claude --- projects/miopen/test/gtest/miopen_cstdint_lp64.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/miopen/test/gtest/miopen_cstdint_lp64.cpp b/projects/miopen/test/gtest/miopen_cstdint_lp64.cpp index 2b8a2cf13926..cc87b4c9e12d 100644 --- a/projects/miopen/test/gtest/miopen_cstdint_lp64.cpp +++ b/projects/miopen/test/gtest/miopen_cstdint_lp64.cpp @@ -49,19 +49,19 @@ typedef signed long long int64_t; #undef MIOPEN_HIP_RUNTIME_COMPILE -TEST(MiopenCstdintLP64, Uint64TypeMatchesSystem) +TEST(CPU_MiopenCstdintLP64_NONE, Uint64TypeMatchesSystem) { EXPECT_TRUE((std::is_same<::uint64_t, std::uint64_t>::value)) << "miopen_cstdint.hpp uint64_t must be the same type as uint64_t"; } -TEST(MiopenCstdintLP64, Int64TypeMatchesSystem) +TEST(CPU_MiopenCstdintLP64_NONE, Int64TypeMatchesSystem) { EXPECT_TRUE((std::is_same<::int64_t, std::int64_t>::value)) << "miopen_cstdint.hpp int64_t must be the same type as int64_t"; } -TEST(MiopenCstdintLP64, SizesAre64Bit) +TEST(CPU_MiopenCstdintLP64_NONE, SizesAre64Bit) { EXPECT_EQ(sizeof(::uint64_t), 8u); EXPECT_EQ(sizeof(::int64_t), 8u); From 01b8cc4a7af158483d347317afd92f78c62d4c6e Mon Sep 17 00:00:00 2001 From: Steven Chung Date: Mon, 28 Sep 2026 10:02:58 -0400 Subject: [PATCH 3/5] =?UTF-8?q?fix(miopen):=20address=20review=20=E2=80=94?= =?UTF-8?q?=20use=20=5F=5FUINT64=5FTYPE=5F=5F,=20SPDX=20header,=20JIT=20te?= =?UTF-8?q?st?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit - Use __UINT64_TYPE__ / __INT64_TYPE__ compiler builtins instead of __LP64__ guards (reviewer suggestion) - Switch test to SPDX license header - Rewrite test to exercise the real comgr JIT path via Handle::AddKernel with inline kernel source including miopen_cstdint.hpp Co-Authored-By: Claude --- .../miopen/src/kernels/miopen_cstdint.hpp | 9 +- .../miopen/test/gtest/miopen_cstdint_lp64.cpp | 87 ++++++++----------- 2 files changed, 37 insertions(+), 59 deletions(-) diff --git a/projects/miopen/src/kernels/miopen_cstdint.hpp b/projects/miopen/src/kernels/miopen_cstdint.hpp index cb1eff85bc31..ed4a4651bc93 100644 --- a/projects/miopen/src/kernels/miopen_cstdint.hpp +++ b/projects/miopen/src/kernels/miopen_cstdint.hpp @@ -33,13 +33,8 @@ typedef unsigned short uint16_t; #if HIP_PACKAGE_VERSION_FLAT >= 6000025000ULL typedef signed int int32_t; typedef unsigned int uint32_t; -#ifdef __LP64__ -typedef unsigned long uint64_t; -typedef signed long int64_t; -#else -typedef unsigned long long uint64_t; -typedef signed long long int64_t; -#endif +typedef __UINT64_TYPE__ uint64_t; +typedef __INT64_TYPE__ int64_t; #endif #else diff --git a/projects/miopen/test/gtest/miopen_cstdint_lp64.cpp b/projects/miopen/test/gtest/miopen_cstdint_lp64.cpp index cc87b4c9e12d..bf4bcb61c3af 100644 --- a/projects/miopen/test/gtest/miopen_cstdint_lp64.cpp +++ b/projects/miopen/test/gtest/miopen_cstdint_lp64.cpp @@ -1,68 +1,51 @@ -/******************************************************************************* - * - * MIT License - * - * Copyright (c) 2026 Advanced Micro Devices, Inc. - * - * 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. - * - *******************************************************************************/ +// SPDX-License-Identifier: MIT +// SPDX-FileCopyrightText: 2026 Advanced Micro Devices, Inc. // Regression test for TheRock#6258 / LP64 typedef conflict in miopen_cstdint.hpp. // // On LP64 platforms (64-bit Linux, macOS), defines uint64_t as // 'unsigned long', but __hip_internal::uint64_t is 'unsigned long long'. // These are distinct C++ types despite being the same width, so a typedef -// redefinition is a hard error when both are visible. +// redefinition is a hard error when both are visible during JIT compilation. +// +// This test exercises the real comgr JIT path via Handle::AddKernel with an +// inline kernel that includes miopen_cstdint.hpp alongside system headers. #include -#include -#include - -#define MIOPEN_HIP_RUNTIME_COMPILE -#define HIP_PACKAGE_VERSION_FLAT 7015026333ULL - -// Provide __hip_internal types as they would appear during HIPRTC compilation -namespace __hip_internal { -typedef unsigned long long uint64_t; -typedef signed long long int64_t; -} // namespace __hip_internal -// Include the header under test — this must not conflict with above -#include "miopen_cstdint.hpp" +#include +#include -#undef MIOPEN_HIP_RUNTIME_COMPILE - -TEST(CPU_MiopenCstdintLP64_NONE, Uint64TypeMatchesSystem) -{ - EXPECT_TRUE((std::is_same<::uint64_t, std::uint64_t>::value)) - << "miopen_cstdint.hpp uint64_t must be the same type as uint64_t"; -} +#include "get_handle.hpp" -TEST(CPU_MiopenCstdintLP64_NONE, Int64TypeMatchesSystem) +static std::string CstdintKernelSource() { - EXPECT_TRUE((std::is_same<::int64_t, std::int64_t>::value)) - << "miopen_cstdint.hpp int64_t must be the same type as int64_t"; + return "#include \"miopen_cstdint.hpp\"\n" + "extern \"C\" {\n" + "__global__ void cstdint_write(uint64_t* data) {\n" + " if(threadIdx.x == 0 && blockIdx.x == 0)\n" + " data[0] = (uint64_t)42;\n" + "}\n" + "}\n"; } -TEST(CPU_MiopenCstdintLP64_NONE, SizesAre64Bit) +TEST(CPU_MiopenCstdintLP64_NONE, JitCompileWithMiopenCstdint) { - EXPECT_EQ(sizeof(::uint64_t), 8u); - EXPECT_EQ(sizeof(::int64_t), 8u); + auto&& h = get_handle(); + + std::vector data_in(1, 0); + auto data_dev = h.Write(data_in); + + h.AddKernel("NoAlgo", + "", + "cstdint_test.cpp", + "cstdint_write", + {1, 1, 1}, + {1, 1, 1}, + "", + 0, + CstdintKernelSource())(data_dev.get()); + + auto data_out = h.Read(data_dev, 1); + EXPECT_EQ(data_out[0], 42u); } From 194b56faec2d7f260a222e2b27204054742ff870 Mon Sep 17 00:00:00 2001 From: Steven Chung Date: Mon, 28 Sep 2026 10:08:03 -0400 Subject: [PATCH 4/5] fix(miopen): include in JIT test to reproduce original conflict The original bug requires both system and miopen_cstdint.hpp to be visible during JIT compilation. Include explicitly in the test kernel source to ensure the typedef conflict would be triggered with the old code, matching the MIOpenNeuron.cpp failure pattern. Co-Authored-By: Claude --- projects/miopen/test/gtest/miopen_cstdint_lp64.cpp | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) diff --git a/projects/miopen/test/gtest/miopen_cstdint_lp64.cpp b/projects/miopen/test/gtest/miopen_cstdint_lp64.cpp index bf4bcb61c3af..fe9568120eb5 100644 --- a/projects/miopen/test/gtest/miopen_cstdint_lp64.cpp +++ b/projects/miopen/test/gtest/miopen_cstdint_lp64.cpp @@ -20,7 +20,14 @@ static std::string CstdintKernelSource() { - return "#include \"miopen_cstdint.hpp\"\n" + // Deliberately include before miopen_cstdint.hpp to reproduce + // the original conflict: system uint64_t (unsigned long) vs the MIOpen + // typedef (previously unsigned long long via __hip_internal). + return "#ifndef MIOPEN_HIP_RUNTIME_COMPILE\n" + "#include \n" + "#endif\n" + "#include \n" + "#include \"miopen_cstdint.hpp\"\n" "extern \"C\" {\n" "__global__ void cstdint_write(uint64_t* data) {\n" " if(threadIdx.x == 0 && blockIdx.x == 0)\n" From 23270dce0eb2ff90ec611c01037fc7c193a5015e Mon Sep 17 00:00:00 2001 From: Steven Chung Date: Mon, 28 Sep 2026 10:16:18 -0400 Subject: [PATCH 5/5] fix(miopen): match repo SPDX license convention Co-Authored-By: Claude --- projects/miopen/test/gtest/miopen_cstdint_lp64.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/miopen/test/gtest/miopen_cstdint_lp64.cpp b/projects/miopen/test/gtest/miopen_cstdint_lp64.cpp index fe9568120eb5..002263333ff8 100644 --- a/projects/miopen/test/gtest/miopen_cstdint_lp64.cpp +++ b/projects/miopen/test/gtest/miopen_cstdint_lp64.cpp @@ -1,5 +1,5 @@ -// SPDX-License-Identifier: MIT -// SPDX-FileCopyrightText: 2026 Advanced Micro Devices, Inc. +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT // Regression test for TheRock#6258 / LP64 typedef conflict in miopen_cstdint.hpp. //