https://github.com/ayokunle321 created https://github.com/llvm/llvm-project/pull/223104
Adds codegen for the following AMDGCN bitop3 builtins: - __builtin_amdgcn_bitop3_b32 - __builtin_amdgcn_bitop3_b16 These are lowered to the corresponding `llvm.amdgcn.bitop3` intrinsic. >From 4316532d326924931f9cb8a1b3d088b20c1b798d Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Sat, 12 Sep 2026 01:48:55 +0200 Subject: [PATCH] add bitop3 builtins --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 9 ++--- .../CodeGenHIP/builtins-amdgcn-gfx1250.hip | 19 +++++++++++ .../CIR/CodeGenHIP/builtins-amdgcn-gfx950.hip | 33 +++++++++++++++++++ 3 files changed, 55 insertions(+), 6 deletions(-) create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx950.hip diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 1c23ea142bf8d..9b47d2389d244 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -988,12 +988,9 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, return mlir::Value{}; } case AMDGPU::BI__builtin_amdgcn_bitop3_b32: - case AMDGPU::BI__builtin_amdgcn_bitop3_b16: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; - } + case AMDGPU::BI__builtin_amdgcn_bitop3_b16: + return emitBuiltinWithOneOverloadedType<4>(expr, "amdgcn.bitop3") + .getValue(); case AMDGPU::BI__builtin_amdgcn_make_buffer_rsrc: { cgm.errorNYI(expr->getSourceRange(), std::string("unimplemented AMDGPU builtin call: ") + diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip index b6e46ba180c34..16f094b15bdfc 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip @@ -12,6 +12,7 @@ // RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s #define __device__ __attribute__((device)) +#define global __attribute__((address_space(1))) //===----------------------------------------------------------------------===// // Test AMDGPU builtins @@ -96,3 +97,21 @@ __device__ void test_tanh_f16(_Float16* out, _Float16 a) { __device__ void test_tanh_bf16(__bf16* out, __bf16 a) { *out = __builtin_amdgcn_tanh_bf16(a); } + +// CIR-LABEL: @_Z15test_bitop3_b32PU3AS1jjjj +// CIR: cir.call_llvm_intrinsic "amdgcn.bitop3" {{.*}} : (!s32i, !s32i, !s32i, !u32i) -> !s32i +// LLVM: define{{.*}} void @_Z15test_bitop3_b32PU3AS1jjjj +// LLVM: call{{.*}} i32 @llvm.amdgcn.bitop3.i32(i32 %{{.+}}, i32 %{{.+}}, i32 %{{.+}}, i32 1) +__device__ void test_bitop3_b32(global unsigned int* out, unsigned int a, + unsigned int b, unsigned int c) { + *out = __builtin_amdgcn_bitop3_b32(a, b, c, 1); +} + +// CIR-LABEL: @_Z15test_bitop3_b16PU3AS1tttt +// CIR: cir.call_llvm_intrinsic "amdgcn.bitop3" {{.*}} : (!s16i, !s16i, !s16i, !u32i) -> !s16i +// LLVM: define{{.*}} void @_Z15test_bitop3_b16PU3AS1tttt +// LLVM: call{{.*}} i16 @llvm.amdgcn.bitop3.i16(i16 %{{.+}}, i16 %{{.+}}, i16 %{{.+}}, i32 1) +__device__ void test_bitop3_b16(global unsigned short* out, unsigned short a, + unsigned short b, unsigned short c) { + *out = __builtin_amdgcn_bitop3_b16(a, b, c, 1); +} diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx950.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx950.hip new file mode 100644 index 0000000000000..4fd0e12d6b043 --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx950.hip @@ -0,0 +1,33 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgpu9.50-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple amdgpu9.50-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgpu9.50-amd-amdhsa -x hip -std=c++11 \ +// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +#define __device__ __attribute__((device)) +#define global __attribute__((address_space(1))) + +// CIR-LABEL: @_Z15test_bitop3_b32PU3AS1jjjj +// CIR: cir.call_llvm_intrinsic "amdgcn.bitop3" {{.*}} : (!s32i, !s32i, !s32i, !u32i) -> !s32i +// LLVM: define{{.*}} void @_Z15test_bitop3_b32PU3AS1jjjj +// LLVM: call{{.*}} i32 @llvm.amdgcn.bitop3.i32(i32 %{{.+}}, i32 %{{.+}}, i32 %{{.+}}, i32 1) +__device__ void test_bitop3_b32(global unsigned int* out, unsigned int a, + unsigned int b, unsigned int c) { + *out = __builtin_amdgcn_bitop3_b32(a, b, c, 1); +} + +// CIR-LABEL: @_Z15test_bitop3_b16PU3AS1tttt +// CIR: cir.call_llvm_intrinsic "amdgcn.bitop3" {{.*}} : (!s16i, !s16i, !s16i, !u32i) -> !s16i +// LLVM: define{{.*}} void @_Z15test_bitop3_b16PU3AS1tttt +// LLVM: call{{.*}} i16 @llvm.amdgcn.bitop3.i16(i16 %{{.+}}, i16 %{{.+}}, i16 %{{.+}}, i32 1) +__device__ void test_bitop3_b16(global unsigned short* out, unsigned short a, + unsigned short b, unsigned short c) { + *out = __builtin_amdgcn_bitop3_b16(a, b, c, 1); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
