Repository navigation
fix(rocRAND): give internal kernels hidden visibility - #13363
Merged
amd-mcieslak merged 2 commits intoOct 10, 2026
Merged
Conversation
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
Pre-commit check failed⛔ pre-commit failed Please run locally:
This repo uses |
JIRA ID : AIPRIMS-247
## Motivation
rocRAND's internal kernels leak into the dynamic symbol table of `librocrand.so`. This exposes implementation details as part of the library's ABI, bloats the symbol table, and risks symbol clashes with other libraries or user code. Of 6,484 dynamic exports, 6,426 were host-side symbols for instantiations of two internal kernel templates.
## Technical Details
- `library/CMakeLists.txt` already sets `CXX_VISIBILITY_PRESET hidden` / `VISIBILITY_INLINES_HIDDEN ON`, and `-fvisibility=hidden` reaches every compile command.
- Added `ROCRAND_INTERNAL_KERNEL` to the internal-only header `library/src/rng/common.hpp`. It expands to `__global__ __attribute__((__visibility__("hidden")))`, which mirrors `ROCPRIM_KERNEL` in rocPRIM's `config.hpp`. It falls back to plain `__global__` for host compilers without GNU attribute support (e.g. nvcc + MSVC).
- Replaced the bare `__global__` on all three internal kernel definitions.
## Test Plan
- Compare dynamic exports before and after: `nm -D --defined-only librocrand.so | c++filt`
- Run the full rocRAND test suite (`ctest`) to confirm the hidden kernels are still found and launched at runtime.
## Test Result
- Dynamic exports: 6,484 before, 58 after. No `trampoline_kernel` / `generate_sobol_kernel` symbols remain.
- `ctest`: 100% tests passed, 0 tests failed out of 53 (gfx950, 183.58 s).
## Submission Checklist
- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
amd-mcieslak
force-pushed
the
users/amd-mcieslak/rocrand-hidden-kernel-visibility
branch
from
October 8, 2026 22:16
7e344fc to
6679289
Compare
|
🎉 All checks passed! This PR is ready for review. |
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #13363 +/- ##
===========================================
+ Coverage 71.86% 72.16% +0.30%
===========================================
Files 2876 2874 -2
Lines 475315 474877 -438
Branches 70237 70143 -94
===========================================
+ Hits 341569 342674 +1105
+ Misses 109960 108400 -1560
- Partials 23786 23803 +17
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
amd-mcieslak
deleted the
users/amd-mcieslak/rocrand-hidden-kernel-visibility
branch
October 10, 2026 17:39
assistant-librarian Bot
pushed a commit
to ROCm/rocRAND
that referenced
this pull request
Oct 10, 2026
fix(rocRAND): give internal kernels hidden visibility
(#13363)
JIRA ID : AIPRIMS-247
## Motivation
rocRAND's internal kernels leak into the dynamic symbol table of
`librocrand.so`. This exposes implementation details as part of the
library's ABI, bloats the symbol table, and risks symbol clashes with
other libraries or user code. Of 6,484 dynamic exports, 6,426 were
host-side symbols for instantiations of two internal kernel templates.
## Technical Details
- `library/CMakeLists.txt` already sets `CXX_VISIBILITY_PRESET hidden` /
`VISIBILITY_INLINES_HIDDEN ON`, and `-fvisibility=hidden` reaches every
compile command.
- Added `ROCRAND_INTERNAL_KERNEL` to the internal-only header
`library/src/rng/common.hpp`. It expands to `__global__
__attribute__((__visibility__("hidden")))`, which mirrors
`ROCPRIM_KERNEL` in rocPRIM's `config.hpp`. It falls back to plain
`__global__` for host compilers without GNU attribute support (e.g. nvcc
+ MSVC).
- Replaced the bare `__global__` on all three internal kernel
definitions.
## Test Plan
- Compare dynamic exports before and after: `nm -D --defined-only
librocrand.so | c++filt`
- Run the full rocRAND test suite (`ctest`) to confirm the hidden
kernels are still found and launched at runtime.
## Test Result
- Dynamic exports: 6,484 before, 58 after. No `trampoline_kernel` /
`generate_sobol_kernel` symbols remain.
- `ctest`: 100% tests passed, 0 tests failed out of 53 (gfx950, 183.58
s).
## Submission Checklist
- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
Contributor
TheRock Submodule Bump ActivityNewest first
|
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
JIRA ID : AIPRIMS-247
Motivation
rocRAND's internal kernels leak into the dynamic symbol table of
librocrand.so. This exposes implementation details as part of the library's ABI, bloats the symbol table, and risks symbol clashes with other libraries or user code. Of 6,484 dynamic exports, 6,426 were host-side symbols for instantiations of two internal kernel templates.Technical Details
library/CMakeLists.txtalready setsCXX_VISIBILITY_PRESET hidden/VISIBILITY_INLINES_HIDDEN ON, and-fvisibility=hiddenreaches every compile command.ROCRAND_INTERNAL_KERNELto the internal-only headerlibrary/src/rng/common.hpp. It expands to__global__ __attribute__((__visibility__("hidden"))), which mirrorsROCPRIM_KERNELin rocPRIM'sconfig.hpp. It falls back to plain__global__for host compilers without GNU attribute support (e.g. nvcc + MSVC).__global__on all three internal kernel definitions.Test Plan
nm -D --defined-only librocrand.so | c++filtctest) to confirm the hidden kernels are still found and launched at runtime.Test Result
trampoline_kernel/generate_sobol_kernelsymbols remain.ctest: 100% tests passed, 0 tests failed out of 53 (gfx950, 183.58 s).Submission Checklist