[clang][AMDGPU] fix the return type for ballot

[ssahasra]

In the builtins declaration, "ULi" is a 32-bit integer on Windows. Use "WUi" instead to ensure a 64-bit integer on all platforms.

[llvmbot]

@llvm/pr-subscribers-backend-amdgpu

@llvm/pr-subscribers-clang

Author: Sameer Sahasrabuddhe (ssahasra)

Changes

In the builtins declaration, "ULi" is a 32-bit integer on Windows. Use "WUi" instead to ensure a 64-bit integer on all platforms.

---

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

2 Files Affected:

• (modified) clang/include/clang/Basic/BuiltinsAMDGPU.def (+2-2)
• (added) clang/test/CodeGenHIP/ballot.cpp (+15)

diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index a19c8bd5f219ec6..8b59b3790d7bc66 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -150,8 +150,8 @@ BUILTIN(__builtin_amdgcn_mqsad_u32_u8, "V4UiWUiUiV4Ui", "nc")
 // Ballot builtins.
 //===----------------------------------------------------------------------===//
 
-TARGET_BUILTIN(__builtin_amdgcn_ballot_w32, "Uib", "nc", "wavefrontsize32")
-TARGET_BUILTIN(__builtin_amdgcn_ballot_w64, "LUib", "nc", "wavefrontsize64")
+TARGET_BUILTIN(__builtin_amdgcn_ballot_w32, "ZUib", "nc", "wavefrontsize32")
+TARGET_BUILTIN(__builtin_amdgcn_ballot_w64, "WUib", "nc", "wavefrontsize64")
 
 // Deprecated intrinsics in favor of __builtin_amdgn_ballot_{w32|w64}
 BUILTIN(__builtin_amdgcn_uicmp, "WUiUiUiIi", "nc")
diff --git a/clang/test/CodeGenHIP/ballot.cpp b/clang/test/CodeGenHIP/ballot.cpp
new file mode 100644
index 000000000000000..5685c5cad30d66a
--- /dev/null
+++ b/clang/test/CodeGenHIP/ballot.cpp
@@ -0,0 +1,15 @@
+// REQUIRES: amdgpu-registered-target
+// 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
+
+// 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);
+}

[arsenm]

We should probably do this for every single builtin

[ssahasra]

I checked now. The tell is whether the builtin uses "L" in its type descriptor. None of them do except ballot.

[arsenm]

Is this testing anything different than the existing builtin test without a windows host component?

[ssahasra]

You're right. The precommit test needed fixing. The updated version uses an aux-triple to specify a Windows host. Now the test exposes the problem even on a Linux build when trying cross-compile to Windows.

Added another note in the test: The same problem does not arise with existing OpenCL tests because the language fully specifies that "long" is 64 bits on all platforms.