https://github.com/ssahasra updated https://github.com/llvm/llvm-project/pull/73906
>From 8ecb6310a4912de50628cf3db5ff8488fa919bb1 Mon Sep 17 00:00:00 2001 From: Sameer Sahasrabuddhe <sameer.sahasrabud...@amd.com> Date: Fri, 1 Dec 2023 14:24:30 +0530 Subject: [PATCH 1/2] [clang][AMDGPU] precommit test for ballot on Windows 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. --- clang/test/CodeGenHIP/ballot.cpp | 27 +++++++++++++++++++++++++++ 1 file changed, 27 insertions(+) create mode 100644 clang/test/CodeGenHIP/ballot.cpp diff --git a/clang/test/CodeGenHIP/ballot.cpp b/clang/test/CodeGenHIP/ballot.cpp new file mode 100644 index 000000000000000..6e1cbbdfc7af170 --- /dev/null +++ b/clang/test/CodeGenHIP/ballot.cpp @@ -0,0 +1,27 @@ +// REQUIRES: amdgpu-registered-target +// XFAIL: * +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -aux-triple x86_64-pc-windows-msvc -target-cpu gfx900 -x hip -emit-llvm -fcuda-is-device -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -aux-triple x86_64-pc-windows-msvc -target-cpu gfx900 -x hip -S -fcuda-is-device -o - %s | FileCheck %s --check-prefix=GFX9 + +// Unlike OpenCL, HIP depends on the C++ interpration of "unsigned long", which +// is 64 bits long on Linux and 32 bits long on Windows. The return type of the +// ballot intrinsic needs to be a 64 bit integer on both platforms. This test +// cross-compiles to Windows to confirm that the return type is indeed 64 bits +// on Windows. + +// 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); +} >From bfcff343a601923da554cafda26568a445fc39b0 Mon Sep 17 00:00:00 2001 From: Sameer Sahasrabuddhe <sameer.sahasrabud...@amd.com> Date: Thu, 30 Nov 2023 12:14:38 +0530 Subject: [PATCH 2/2] [clang][AMDGPU] fix the return type for ballot In the builtins declaration, "ULi" is a 32-bit integer on Windows. Use "WUi" instead to ensure a 64-bit integer on all platforms. --- clang/include/clang/Basic/BuiltinsAMDGPU.def | 4 ++-- clang/test/CodeGenHIP/ballot.cpp | 6 ------ 2 files changed, 2 insertions(+), 8 deletions(-) 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 index 6e1cbbdfc7af170..a1c23e2136c7153 100644 --- a/clang/test/CodeGenHIP/ballot.cpp +++ b/clang/test/CodeGenHIP/ballot.cpp @@ -1,5 +1,4 @@ // REQUIRES: amdgpu-registered-target -// XFAIL: * // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -aux-triple x86_64-pc-windows-msvc -target-cpu gfx900 -x hip -emit-llvm -fcuda-is-device -o - %s | FileCheck %s // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -aux-triple x86_64-pc-windows-msvc -target-cpu gfx900 -x hip -S -fcuda-is-device -o - %s | FileCheck %s --check-prefix=GFX9 @@ -9,11 +8,6 @@ // cross-compiles to Windows to confirm that the return type is indeed 64 bits // on Windows. -// 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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits