Skip to content

[AMDGPU] Add missing __builtin_amdgcn_wavefrontsize builtin #80741

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Feb 5, 2024

Conversation

jhuber6
Copy link
Contributor

@jhuber6 jhuber6 commented Feb 5, 2024

Summary:
The backend supports the wavefrontsize intrinsic, and suggests that it
is tied to a corresponding clang builtin, but it is not actually
present. This simply adds it in so it can be used from clang. This
attribute likely isn't the best to rely on, but for the libc use-case
we will need to detect a struct's differing size in a way that will
depend on the wavefront size.

Summary:
The backend supports the wavefrontsize intrinsic, and suggests that it
is tied to a corresponding clang builtin, but it is not actually
present. This simply adds it in so it can be used from clang. This
attribute likely isn't the best to rely on, but for the `libc` use-case
we will need to detect a struct's differing size in a way that will
depend on the wavefront size.
@llvmbot llvmbot added clang Clang issues not falling into any other category backend:AMDGPU clang:frontend Language frontend issues, e.g. anything involving "Sema" labels Feb 5, 2024
@llvmbot
Copy link
Member

llvmbot commented Feb 5, 2024

@llvm/pr-subscribers-clang

@llvm/pr-subscribers-backend-amdgpu

Author: Joseph Huber (jhuber6)

Changes

Summary:
The backend supports the wavefrontsize intrinsic, and suggests that it
is tied to a corresponding clang builtin, but it is not actually
present. This simply adds it in so it can be used from clang. This
attribute likely isn't the best to rely on, but for the libc use-case
we will need to detect a struct's differing size in a way that will
depend on the wavefront size.


Full diff: https://github.com/llvm/llvm-project/pull/80741.diff

2 Files Affected:

  • (modified) clang/include/clang/Basic/BuiltinsAMDGPU.def (+1)
  • (modified) clang/test/CodeGenOpenCL/builtins-amdgcn.cl (+7)
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index 5f8001e61a028..213311b96df74 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -69,6 +69,7 @@ BUILTIN(__builtin_amdgcn_s_dcache_inv, "v", "n")
 BUILTIN(__builtin_amdgcn_buffer_wbinvl1, "v", "n")
 BUILTIN(__builtin_amdgcn_fence, "vUicC*", "n")
 BUILTIN(__builtin_amdgcn_groupstaticsize, "Ui", "n")
+BUILTIN(__builtin_amdgcn_wavefrontsize, "Ui", "nc")
 
 BUILTIN(__builtin_amdgcn_atomic_inc32, "UZiUZiD*UZiUicC*", "n")
 BUILTIN(__builtin_amdgcn_atomic_inc64, "UWiUWiD*UWiUicC*", "n")
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
index 8d9e4e018b12e..7d9010ee9067d 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -832,6 +832,13 @@ void test_atomic_inc_dec(local uint *lptr, global uint *gptr, uint val) {
   res = __builtin_amdgcn_atomic_dec32((volatile global uint*)gptr, val, __ATOMIC_SEQ_CST, "");
 }
 
+// CHECK-LABEL test_wavefrontsize(
+unsigned test_wavefrontsize() {
+
+  // CHECK: call i32 @llvm.amdgcn.wavefrontsize()
+  return __builtin_amdgcn_wavefrontsize();
+}
+
 // CHECK-DAG: [[$WI_RANGE]] = !{i32 0, i32 1024}
 // CHECK-DAG: [[$WS_RANGE]] = !{i16 1, i16 1025}
 // CHECK-DAG: attributes #[[$NOUNWIND_READONLY]] = { convergent mustprogress nocallback nofree nounwind willreturn memory(none) }

@@ -832,6 +832,13 @@ void test_atomic_inc_dec(local uint *lptr, global uint *gptr, uint val) {
res = __builtin_amdgcn_atomic_dec32((volatile global uint*)gptr, val, __ATOMIC_SEQ_CST, "");
}

// CHECK-LABEL test_wavefrontsize(
unsigned test_wavefrontsize() {
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Missing check for the test.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Which check? This just makes sure that we generate the appropriate LLVM intrinsic below.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Ugh, it's inside the body. Unusual, but test above is he same.

@@ -832,6 +832,13 @@ void test_atomic_inc_dec(local uint *lptr, global uint *gptr, uint val) {
res = __builtin_amdgcn_atomic_dec32((volatile global uint*)gptr, val, __ATOMIC_SEQ_CST, "");
}

// CHECK-LABEL test_wavefrontsize(
unsigned test_wavefrontsize() {
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Ugh, it's inside the body. Unusual, but test above is he same.

@jhuber6 jhuber6 merged commit dbed898 into llvm:main Feb 5, 2024
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
backend:AMDGPU clang:frontend Language frontend issues, e.g. anything involving "Sema" clang Clang issues not falling into any other category
Projects
None yet
Development

Successfully merging this pull request may close these issues.

3 participants