https://github.com/ayokunle321 updated https://github.com/llvm/llvm-project/pull/223101
>From e62781adc9a7dfb7c2384140b0b96389fd187ed4 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Wed, 9 Sep 2026 13:07:19 +0200 Subject: [PATCH 1/2] add codegen --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 16 ++++--- .../builtins-amdgcn-gfx1250-cluster-load.hip | 45 +++++++++++++++++++ 2 files changed, 55 insertions(+), 6 deletions(-) create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250-cluster-load.hip diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 1c23ea142bf8d..405d2998355ea 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -507,13 +507,17 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, return mlir::Value{}; } case AMDGPU::BI__builtin_amdgcn_cluster_load_b32: + return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.cluster.load.b32", + convertType(expr->getType())) + .getValue(); case AMDGPU::BI__builtin_amdgcn_cluster_load_b64: - case AMDGPU::BI__builtin_amdgcn_cluster_load_b128: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; - } + return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.cluster.load.b64", + convertType(expr->getType())) + .getValue(); + case AMDGPU::BI__builtin_amdgcn_cluster_load_b128: + return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.cluster.load.b128", + convertType(expr->getType())) + .getValue(); case AMDGPU::BI__builtin_amdgcn_load_to_lds: { cgm.errorNYI(expr->getSourceRange(), std::string("unimplemented AMDGPU builtin call: ") + diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250-cluster-load.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250-cluster-load.hip new file mode 100644 index 0000000000000..8ac92d47b9846 --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250-cluster-load.hip @@ -0,0 +1,45 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgpu12.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 amdgpu12.50-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -fcuda-is-device -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s + +// RUN: %clang_cc1 -triple amdgpu12.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 + +// These intrinsics are overloaded on their result type, not on an argument, so +// the result type is passed to the emission helper explicitly. + +#define __device__ __attribute__((device)) +#define global __attribute__((address_space(1))) + +typedef int v2i __attribute__((ext_vector_type(2))); +typedef int v4i __attribute__((ext_vector_type(4))); + +// CIR-LABEL: @_Z28test_amdgcn_cluster_load_b32PU3AS1ii +// CIR: cir.call_llvm_intrinsic "amdgcn.cluster.load.b32" {{.*}} : (!cir.ptr<!s32i, target_address_space(1)>, !s32i, !s32i) -> !s32i +// LLVM: define{{.*}} i32 @_Z28test_amdgcn_cluster_load_b32PU3AS1ii +// LLVM: call{{.*}} i32 @llvm.amdgcn.cluster.load.b32.i32(ptr addrspace(1) %{{.+}}, i32 10, i32 %{{.+}}) +__device__ int test_amdgcn_cluster_load_b32(global int* inptr, int mask) { + return __builtin_amdgcn_cluster_load_b32(inptr, 10, mask); +} + +// CIR-LABEL: @_Z28test_amdgcn_cluster_load_b64PU3AS1Dv2_ii +// CIR: cir.call_llvm_intrinsic "amdgcn.cluster.load.b64" {{.*}} : (!cir.ptr<!cir.vector<2 x !s32i>, target_address_space(1)>, !s32i, !s32i) -> !cir.vector<2 x !s32i> +// LLVM: define{{.*}} <2 x i32> @_Z28test_amdgcn_cluster_load_b64PU3AS1Dv2_ii +// LLVM: call{{.*}} <2 x i32> @llvm.amdgcn.cluster.load.b64.v2i32(ptr addrspace(1) %{{.+}}, i32 22, i32 %{{.+}}) +__device__ v2i test_amdgcn_cluster_load_b64(global v2i* inptr, int mask) { + return __builtin_amdgcn_cluster_load_b64(inptr, 22, mask); +} + +// CIR-LABEL: @_Z29test_amdgcn_cluster_load_b128PU3AS1Dv4_ii +// CIR: cir.call_llvm_intrinsic "amdgcn.cluster.load.b128" {{.*}} : (!cir.ptr<!cir.vector<4 x !s32i>, target_address_space(1)>, !s32i, !s32i) -> !cir.vector<4 x !s32i> +// LLVM: define{{.*}} <4 x i32> @_Z29test_amdgcn_cluster_load_b128PU3AS1Dv4_ii +// LLVM: call{{.*}} <4 x i32> @llvm.amdgcn.cluster.load.b128.v4i32(ptr addrspace(1) %{{.+}}, i32 27, i32 %{{.+}}) +__device__ v4i test_amdgcn_cluster_load_b128(global v4i* inptr, int mask) { + return __builtin_amdgcn_cluster_load_b128(inptr, 27, mask); +} >From 81966c3bedce6010c5e8b90c2761524d0212c53f Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Sat, 12 Sep 2026 01:31:29 +0200 Subject: [PATCH 2/2] add fmed3 builtins --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 8 +--- .../CIR/CodeGenHIP/builtins-amdgcn-gfx9.hip | 38 +++++++++++++++++++ clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip | 8 ++++ 3 files changed, 48 insertions(+), 6 deletions(-) create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx9.hip diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 405d2998355ea..a6b8dc0dd1a88 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -441,12 +441,8 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, convertType(expr->getType())) .getValue(); case AMDGPU::BI__builtin_amdgcn_fmed3f: - case AMDGPU::BI__builtin_amdgcn_fmed3h: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; - } + case AMDGPU::BI__builtin_amdgcn_fmed3h: + return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.fmed3").getValue(); case AMDGPU::BI__builtin_amdgcn_ds_append: case AMDGPU::BI__builtin_amdgcn_ds_consume: { cgm.errorNYI(expr->getSourceRange(), diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx9.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx9.hip new file mode 100644 index 0000000000000..a082b2553ba1c --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx9.hip @@ -0,0 +1,38 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgpu9.00-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 amdgpu10.10-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.00-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 amdgpu10.10-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.00-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 + +// RUN: %clang_cc1 -triple amdgpu10.10-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 + +// __builtin_amdgcn_fmed3h requires gfx9-insts, so it lives here rather than in +// builtins-amdgcn.hip, mirroring builtins-amdgcn-gfx9.cl. + +#define __device__ __attribute__((device)) + +// CIR-LABEL: @_Z14test_fmed3_f16PDF16_DF16_DF16_DF16_ +// CIR: cir.call_llvm_intrinsic "amdgcn.fmed3" {{.*}} : (!cir.f16, !cir.f16, !cir.f16) -> !cir.f16 +// LLVM: define{{.*}} void @_Z14test_fmed3_f16PDF16_DF16_DF16_DF16_ +// LLVM: call{{.*}} half @llvm.amdgcn.fmed3.f16(half %{{.+}}, half %{{.+}}, half %{{.+}}) +__device__ void test_fmed3_f16(_Float16* out, _Float16 a, _Float16 b, + _Float16 c) { + *out = __builtin_amdgcn_fmed3h(a, b, c); +} diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip index 9bbfb6747cbf8..41a0f9777b7fb 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip @@ -331,3 +331,11 @@ __device__ void test_class_f32(bool* out, float a, int b) { __device__ void test_class_f64(bool* out, double a, int b) { *out = __builtin_amdgcn_class(a, b); } + +// CIR-LABEL: @_Z14test_fmed3_f32Pffff +// CIR: cir.call_llvm_intrinsic "amdgcn.fmed3" {{.*}} : (!cir.float, !cir.float, !cir.float) -> !cir.float +// LLVM: define{{.*}} void @_Z14test_fmed3_f32Pffff +// LLVM: call{{.*}} float @llvm.amdgcn.fmed3.f32(float %{{.+}}, float %{{.+}}, float %{{.+}}) +__device__ void test_fmed3_f32(float* out, float a, float b, float c) { + *out = __builtin_amdgcn_fmed3f(a, b, c); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
