-
Notifications
You must be signed in to change notification settings - Fork 10.8k
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
[clang][AMDGPU] precommit test for ballot on Windows #73920
Conversation
@llvm/pr-subscribers-clang Author: Sameer Sahasrabuddhe (ssahasra) ChangesThe Clang declaration of the wave-64 builtin uses "UL" as the return type, which is interpreted as a 32-bit unsigned integer on Windows. This emits an incorrect LLVM declaration with i32 return type instead of i64. The clang declaration needs to be fixed to use "WU" instead. Full diff: https://github.com/llvm/llvm-project/pull/73920.diff 1 Files Affected:
diff --git a/clang/test/CodeGenHIP/ballot.cpp b/clang/test/CodeGenHIP/ballot.cpp
new file mode 100644
index 000000000000000..823b1583a42fc02
--- /dev/null
+++ b/clang/test/CodeGenHIP/ballot.cpp
@@ -0,0 +1,21 @@
+// REQUIRES: amdgpu-registered-target
+// XFAIL: system-windows
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx900 -x hip -emit-llvm -fcuda-is-device -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx900 -x hip -S -fcuda-is-device -o - %s | FileCheck %s --check-prefix=GFX9
+
+// FIXME: The Clang declaration of the wave-64 builtin uses "UL" as the return
+// type, which is interpreted as a 32-bit unsigned integer on Windows. This
+// emits an incorrect LLVM declaration with i32 return type instead of i64. The
+// clang declaration needs to be fixed to use "WU" instead.
+
+// CHECK-LABEL: @_Z3fooi
+// CHECK: call i64 @llvm.amdgcn.ballot.i64
+
+// GFX9-LABEL: _Z3fooi:
+// GFX9: v_cmp_ne_u32_e64
+
+#define __device__ __attribute__((device))
+
+__device__ unsigned long long foo(int p) {
+ return __builtin_amdgcn_ballot_w64(p);
+}
|
d5945a3
to
6b87550
Compare
The Clang declaration of the wave-64 builtin uses "UL" as the return type, which is interpreted as a 32-bit unsigned integer on Windows. This emits an incorrect LLVM declaration with i32 return type instead of i64. The clang declaration needs to be fixed to use "WU" instead.
6b87550
to
8ecb631
Compare
// CHECK: call i64 @llvm.amdgcn.ballot.i64 | ||
|
||
// GFX9-LABEL: _Z3fooi: | ||
// GFX9: v_cmp_ne_u32_e64 |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Is it really needed to test codegen here?
Nice find, BTW.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Yeah, we should check codegen because ultimately that was the symptom ... llc crashed when running hipcc on Windows.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
LGTM, thanks!
The Clang declaration of the wave-64 builtin uses "UL" as the return type, which is interpreted as a 32-bit unsigned integer on Windows. This emits an incorrect LLVM declaration with i32 return type instead of i64. The clang declaration needs to be fixed to use "WU" instead.