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/3] 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 1c23ea142bf8dc..405d2998355ea5 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 00000000000000..8ac92d47b98464 --- /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/3] 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 405d2998355ea5..a6b8dc0dd1a88c 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 00000000000000..a082b2553ba1c1 --- /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 9bbfb6747cbf8c..41a0f9777b7fb7 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); +} >From 9bd54cc947f7d5b272c3998e050728e217a4eaa5 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Sat, 12 Sep 2026 03:20:57 +0200 Subject: [PATCH 3/3] remove cluster load --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 16 +++---- .../builtins-amdgcn-gfx1250-cluster-load.hip | 45 ------------------- 2 files changed, 6 insertions(+), 55 deletions(-) delete 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 a6b8dc0dd1a88c..017024a681a1cb 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -503,17 +503,13 @@ 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: - 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_cluster_load_b128: { + cgm.errorNYI(expr->getSourceRange(), + std::string("unimplemented AMDGPU builtin call: ") + + getContext().BuiltinInfo.getName(builtinId)); + return mlir::Value{}; + } 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 deleted file mode 100644 index 8ac92d47b98464..00000000000000 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250-cluster-load.hip +++ /dev/null @@ -1,45 +0,0 @@ -// 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); -} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
