Skip to content

Conversation

@kevinjoseph1995
Copy link

@kevinjoseph1995 kevinjoseph1995 commented Oct 18, 2025

Summary:

Usage of __AMDGCN_WAVEFRONT_SIZE is a compilation error when using amdclang++ shipped with ROCm 7.x.x. For versions >= 7 we delegate to rocprim to get the right constexpr value that can be used in device code. We continue using __AMDGCN_WAVEFRONT_SIZE for older ROCm versions to maintain compatibility.

Note that the usage of kWarpSize in host code is still non-sensical however the current usage of this variable seems to always be in device functions.

See https://rocm.docs.amd.com/en/latest/about/release-notes.html#amdgpu-wavefront-size-compiler-macro-deprecation for deprecation notice.

…in ROCm 7

Summary:
Usage of __AMDGCN_WAVEFRONT_SIZE is a compilation error when using amdclang++ shipped with ROCm 7.x.x.
For versions >= 7 we delegate to rocprim to get the right constexpr value that can be used in __device__
code. We continue using __AMDGCN_WAVEFRONT_SIZE for older ROCm versions to maintain compatibility.

Note that the usage of kWarpSize in host code is still non-sensical however
the current usage of this variable seems to always be in __device__ functions.

See https://rocm.docs.amd.com/en/latest/about/release-notes.html#amdgpu-wavefront-size-compiler-macro-deprecation
for deprecation notice.
@meta-cla
Copy link

meta-cla bot commented Oct 18, 2025

Hi @kevinjoseph1995!

Thank you for your pull request and welcome to our community.

Action Required

In order to merge any pull request (code, docs, etc.), we require contributors to sign our Contributor License Agreement, and we don't seem to have one on file for you.

Process

In order for us to review and merge your suggested changes, please sign at https://code.facebook.com/cla. If you are contributing on behalf of someone else (eg your employer), the individual CLA may not be sufficient and your employer may need to sign the corporate CLA.

Once the CLA is signed, our tooling will perform checks and validations. Afterwards, the pull request will be tagged with CLA signed. The tagging process may take up to 1 hour after signing. Please give it that time before contacting us about it.

If you have received this in error or have any questions, please contact us at [email protected]. Thanks!

@meta-cla
Copy link

meta-cla bot commented Oct 18, 2025

Thank you for signing our Contributor License Agreement. We can now accept your code for this (and any) Meta Open Source project. Thanks!

@alibeklfc alibeklfc added the GPU label Oct 20, 2025
@mnorris11
Copy link

Thanks for this contribution @kevinjoseph1995!
Right now, our ROCm runner in the CI is not running. We are working on getting it back up and running again. After it is, we can run this in the CI and then get it merged.

@kevinjoseph1995
Copy link
Author

Got it, thanks for letting me know! I’ll wait for the ROCm runner to be back online.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants