diff --git a/projects/miopen/src/kernels/miopen_cstdint.hpp b/projects/miopen/src/kernels/miopen_cstdint.hpp index fa8d68f5fa58..ed4a4651bc93 100644 --- a/projects/miopen/src/kernels/miopen_cstdint.hpp +++ b/projects/miopen/src/kernels/miopen_cstdint.hpp @@ -33,8 +33,8 @@ 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; +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 new file mode 100644 index 000000000000..002263333ff8 --- /dev/null +++ b/projects/miopen/test/gtest/miopen_cstdint_lp64.cpp @@ -0,0 +1,58 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +// 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 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 + +#include "get_handle.hpp" + +static std::string CstdintKernelSource() +{ + // 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" + " data[0] = (uint64_t)42;\n" + "}\n" + "}\n"; +} + +TEST(CPU_MiopenCstdintLP64_NONE, JitCompileWithMiopenCstdint) +{ + 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); +}