[schung-amd]

Motivation

miopen_cstdint.hpp causes a hard compilation error on LP64 Linux (all 64-bit Linux and macOS) when system <stdint.h> is also visible during MIOpen's runtime kernel compilation (comgr/HIPRTC JIT path).

The header typedefs uint64_t as __hip_internal::uint64_t (unsigned long long), but on LP64 platforms <stdint.h> defines uint64_t as unsigned long. These are distinct C++ types despite being the same width (8 bytes), so the compiler rejects the redefinition:

error: typedef redefinition with different types
  ('__hip_internal::uint64_t' (aka 'unsigned long long') vs '__uint64_t' (aka 'unsigned long'))

This affects any GPU target where MIOpen JIT-compiles kernels at runtime (i.e. when pre-compiled kernel databases are missing). It is most commonly hit on consumer GPUs like gfx1032 that lack .fdb.txt entries, but can affect any target under the right conditions. 29 MIOpen kernel files include this header.

Reported in ROCm/TheRock#6258

Technical Details

On LP64 platforms (64-bit Linux, macOS):

• unsigned long = 64-bit, used by system <stdint.h> for uint64_t
• unsigned long long = 64-bit, used by __hip_internal::uint64_t
• These are distinct C++ types despite identical size and alignment

On LLP64 platforms (Windows):

• unsigned long = 32-bit
• unsigned long long = 64-bit, matches system uint64_t
• No conflict exists

The fix uses __LP64__ (defined by GCC/Clang on LP64 platforms) to select the correct underlying type for each platform, matching what the system <stdint.h> provides. This is consistent with the existing WIN32 guard in test/gtest/CMakeLists.txt (line 341) which already passes -U__LP64__ on Windows builds.

Reproducer

Minimal reproduction — compile with any HIP-capable clang:

#include <stdint.h>                         // system: uint64_t = unsigned long
typedef unsigned long long uint64_t;        // miopen: uint64_t = unsigned long long
// error: typedef redefinition with different types

This is what happens when comgr compiles a MIOpen kernel that includes both system headers and miopen_cstdint.hpp. The reporter hit it via the MIOpenNeuron.cpp kernel on gfx1032, but any of the 29 kernel files that include miopen_cstdint.hpp can trigger it.

Other scenarios that can trigger this:

• Any GPU target missing pre-compiled kernel DB entries (.fdb.txt), forcing JIT
• Custom MIOpen builds where the kernel DB is not fully populated
• New GPU targets added before their kernel DBs are generated

Issue Tracking

GitHub issue: ROCm/TheRock#6258

Test Plan

Added test/gtest/miopen_cstdint_lp64.cpp — a compile-time regression test that includes both <cstdint> and miopen_cstdint.hpp with the MIOPEN_HIP_RUNTIME_COMPILE path active, and verifies via std::is_same that the resulting uint64_t/int64_t types match the system definitions. The test is auto-discovered by the existing file(GLOB TESTS *.cpp) in the gtest CMakeLists.txt.

Test Result

Verified locally on Ubuntu 24.04 (LP64), ROCm 10.0 (amdclang 23.0.0git), gfx1100:

• Before fix: typedef redefinition with different types error
• After fix: compiles cleanly, types match system <stdint.h>

🤖 Generated with Claude Code

[therock-pr-bot]

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

[therock-pr-bot]

Pre-commit check failed

⛔ pre-commit failed

Please run locally:

• python -m pip install pre-commit
• pre-commit install
• pre-commit run --all-files --show-diff-on-failure

This repo uses .pre-commit-config.yaml.

[therock-pr-bot]

🎉 All checks passed! This PR is ready for review.

[averinevg]

Thanks for the fix. Left some comments.

[averinevg]

LGTM

[assistant-librarian]

TheRock Submodule Bump Activity

Newest first

• 2026-10-01 — Included in TheRock via ROCm/TheRock#8607.