https://github.com/ayokunle321 updated https://github.com/llvm/llvm-project/pull/197526
>From 1d5097cb8f88e4a41d14807bb30c27ed718688b9 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Wed, 13 May 2026 14:23:00 -0400 Subject: [PATCH 1/6] add codegen support for permlane builtin --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 27 ++++- .../CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip | 58 ++++++++++ .../CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip | 106 ++++++++++++++++++ 3 files changed, 186 insertions(+), 5 deletions(-) create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 0a7ba0c194400..0b5d1229451a5 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -211,12 +211,29 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, return mlir::Value{}; } case AMDGPU::BI__builtin_amdgcn_permlane16: - case AMDGPU::BI__builtin_amdgcn_permlanex16: + case AMDGPU::BI__builtin_amdgcn_permlanex16: { + llvm::StringRef intrinsicName = + builtinId == AMDGPU::BI__builtin_amdgcn_permlane16 + ? "amdgcn.permlane16" + : "amdgcn.permlanex16"; + mlir::Value src0 = emitScalarExpr(expr->getArg(0)); + mlir::Value src1 = emitScalarExpr(expr->getArg(1)); + mlir::Value src2 = emitScalarExpr(expr->getArg(2)); + mlir::Value src3 = emitScalarExpr(expr->getArg(3)); + mlir::Value src4 = emitScalarExpr(expr->getArg(4)); + mlir::Value src5 = emitScalarExpr(expr->getArg(5)); + mlir::Value result = + cir::LLVMIntrinsicCallOp::create(builder, getLoc(expr->getExprLoc()), + builder.getStringAttr(intrinsicName), src0.getType(), + {src0, src1, src2, src3, src4, src5}) + .getResult(); + return result; + } case AMDGPU::BI__builtin_amdgcn_permlane64: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + mlir::Value src = emitScalarExpr(expr->getArg(0)); + return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), + "amdgcn.permlane64", src.getType(), + mlir::ValueRange{src}); } case AMDGPU::BI__builtin_amdgcn_readlane: return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.readlane") diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip new file mode 100644 index 0000000000000..bdb3018d1f4d7 --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip @@ -0,0 +1,58 @@ +#include "../CodeGenCUDA/Inputs/cuda.h" + +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1010 -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1011 -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1012 -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1010 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1011 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \ +// RUN: -target-cpu gfx1012 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \ +// RUN: -target-cpu gfx1010 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \ +// RUN: -target-cpu gfx1011 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \ +// RUN: -target-cpu gfx1012 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +//===----------------------------------------------------------------------===// +// Test AMDGPU builtins +//===----------------------------------------------------------------------===// + +// CIR-LABEL: @_Z15test_permlane16Pjjjjj +// CIR: cir.call_llvm_intrinsic "amdgcn.permlane16" {{.*}} : (!u32i, !u32i, !u32i, !u32i, !cir.bool, !cir.bool) -> !u32i +// LLVM: define{{.*}} void @_Z15test_permlane16Pjjjjj +// LLVM: call i32 @llvm.amdgcn.permlane16.i32(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i1 false, i1 false) +__device__ void test_permlane16(unsigned int* out, unsigned int a, unsigned int b, unsigned int c, unsigned int d) { + *out = __builtin_amdgcn_permlane16(a, b, c, d, 0, 0); +} + +// CIR-LABEL: @_Z16test_permlanex16Pjjjjj +// CIR: cir.call_llvm_intrinsic "amdgcn.permlanex16" {{.*}} : (!u32i, !u32i, !u32i, !u32i, !cir.bool, !cir.bool) -> !u32i +// LLVM: define{{.*}} void @_Z16test_permlanex16Pjjjjj +// LLVM: call i32 @llvm.amdgcn.permlanex16.i32(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i1 false, i1 false) +__device__ void test_permlanex16(unsigned int* out, unsigned int a, unsigned int b, unsigned int c, unsigned int d) { + *out = __builtin_amdgcn_permlanex16(a, b, c, d, 0, 0); +} diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip new file mode 100644 index 0000000000000..164ef5a8e2e95 --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip @@ -0,0 +1,106 @@ +#include "../CodeGenCUDA/Inputs/cuda.h" + +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1100 -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1101 -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1102 -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1103 -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1150 -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1151 -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1152 -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1100 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1101 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1102 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1103 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1150 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1151 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1152 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1153 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \ +// RUN: -target-cpu gfx1100 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \ +// RUN: -target-cpu gfx1101 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \ +// RUN: -target-cpu gfx1102 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \ +// RUN: -target-cpu gfx1103 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \ +// RUN: -target-cpu gfx1150 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \ +// RUN: -target-cpu gfx1151 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \ +// RUN: -target-cpu gfx1152 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \ +// RUN: -target-cpu gfx1153 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +//===----------------------------------------------------------------------===// +// Test AMDGPU builtins +//===----------------------------------------------------------------------===// + +// CIR-LABEL: @_Z15test_permlane64Pjj +// CIR: cir.call_llvm_intrinsic "amdgcn.permlane64" {{.*}} : (!u32i) -> !u32i +// LLVM: define{{.*}} void @_Z15test_permlane64Pjj +// LLVM: call i32 @llvm.amdgcn.permlane64.i32(i32 %{{.*}}) +__device__ void test_permlane64(unsigned int* out, unsigned int a) { + *out = __builtin_amdgcn_permlane64(a); +} >From 70e208aa3c834c3b8211996e2efa84ab23cafb6e Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Wed, 13 May 2026 14:39:35 -0400 Subject: [PATCH 2/6] switch out create call --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 12 ++++++++---- 1 file changed, 8 insertions(+), 4 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 0b5d1229451a5..e59d653845468 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -224,16 +224,20 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, mlir::Value src5 = emitScalarExpr(expr->getArg(5)); mlir::Value result = cir::LLVMIntrinsicCallOp::create(builder, getLoc(expr->getExprLoc()), - builder.getStringAttr(intrinsicName), src0.getType(), + builder.getStringAttr(intrinsicName), + src0.getType(), {src0, src1, src2, src3, src4, src5}) .getResult(); return result; } case AMDGPU::BI__builtin_amdgcn_permlane64: { mlir::Value src = emitScalarExpr(expr->getArg(0)); - return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), - "amdgcn.permlane64", src.getType(), - mlir::ValueRange{src}); + mlir::Value result = + cir::LLVMIntrinsicCallOp::create( + builder, getLoc(expr->getExprLoc()), + builder.getStringAttr("amdgcn.permlane64"), src.getType(), {src}) + .getResult(); + return result; } case AMDGPU::BI__builtin_amdgcn_readlane: return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.readlane") >From d5fdeb01eddde084eed2f7e6fa510a8dcfc599c1 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Thu, 14 May 2026 16:13:33 -0400 Subject: [PATCH 3/6] switch out create call --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 19 ++++++------------- 1 file changed, 6 insertions(+), 13 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index e59d653845468..ee7980035d16b 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -222,22 +222,15 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, mlir::Value src3 = emitScalarExpr(expr->getArg(3)); mlir::Value src4 = emitScalarExpr(expr->getArg(4)); mlir::Value src5 = emitScalarExpr(expr->getArg(5)); - mlir::Value result = - cir::LLVMIntrinsicCallOp::create(builder, getLoc(expr->getExprLoc()), - builder.getStringAttr(intrinsicName), - src0.getType(), - {src0, src1, src2, src3, src4, src5}) - .getResult(); - return result; + return builder.emitIntrinsicCallOp( + getLoc(expr->getExprLoc()), intrinsicName, src0.getType(), + mlir::ValueRange{src0, src1, src2, src3, src4, src5}); } case AMDGPU::BI__builtin_amdgcn_permlane64: { mlir::Value src = emitScalarExpr(expr->getArg(0)); - mlir::Value result = - cir::LLVMIntrinsicCallOp::create( - builder, getLoc(expr->getExprLoc()), - builder.getStringAttr("amdgcn.permlane64"), src.getType(), {src}) - .getResult(); - return result; + return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), + "amdgcn.permlane64", src.getType(), + mlir::ValueRange{src}); } case AMDGPU::BI__builtin_amdgcn_readlane: return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.readlane") >From 9adf6a085a4a22d92b81fb5f0bf8ed9707e21d11 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Fri, 10 Jul 2026 15:28:36 -0400 Subject: [PATCH 4/6] switch manual codegen for function --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 16 +++------------- .../CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip | 4 ++++ 2 files changed, 7 insertions(+), 13 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index ee7980035d16b..2479e313fae49 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -216,21 +216,11 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, builtinId == AMDGPU::BI__builtin_amdgcn_permlane16 ? "amdgcn.permlane16" : "amdgcn.permlanex16"; - mlir::Value src0 = emitScalarExpr(expr->getArg(0)); - mlir::Value src1 = emitScalarExpr(expr->getArg(1)); - mlir::Value src2 = emitScalarExpr(expr->getArg(2)); - mlir::Value src3 = emitScalarExpr(expr->getArg(3)); - mlir::Value src4 = emitScalarExpr(expr->getArg(4)); - mlir::Value src5 = emitScalarExpr(expr->getArg(5)); - return builder.emitIntrinsicCallOp( - getLoc(expr->getExprLoc()), intrinsicName, src0.getType(), - mlir::ValueRange{src0, src1, src2, src3, src4, src5}); + return emitBuiltinWithOneOverloadedType<6>(expr, intrinsicName).getValue(); } case AMDGPU::BI__builtin_amdgcn_permlane64: { - mlir::Value src = emitScalarExpr(expr->getArg(0)); - return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), - "amdgcn.permlane64", src.getType(), - mlir::ValueRange{src}); + return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.permlane64") + .getValue(); } case AMDGPU::BI__builtin_amdgcn_readlane: return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.readlane") diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip index 164ef5a8e2e95..304a40b06e284 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip @@ -29,6 +29,10 @@ // RUN: -target-cpu gfx1152 -fcuda-is-device -emit-cir %s -o %t.cir // RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1153 -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ // RUN: -target-cpu gfx1100 -fcuda-is-device -emit-llvm %s -o %t.ll // RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s >From 6509b90d09d24e38ffeea9f19e720ffaf4415f3d Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Fri, 10 Jul 2026 15:30:34 -0400 Subject: [PATCH 5/6] switch header for macro in test --- clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip | 2 ++ clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip | 2 ++ 2 files changed, 4 insertions(+) diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip index bdb3018d1f4d7..aa10cd5b299b7 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip @@ -37,6 +37,8 @@ // RUN: -target-cpu gfx1012 -fcuda-is-device -emit-llvm %s -o %t.ll // RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s +#define __device__ __attribute__((device)) + //===----------------------------------------------------------------------===// // Test AMDGPU builtins //===----------------------------------------------------------------------===// diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip index 304a40b06e284..90fdc2f0d9ae1 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip @@ -97,6 +97,8 @@ // RUN: -target-cpu gfx1153 -fcuda-is-device -emit-llvm %s -o %t.ll // RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s +#define __device__ __attribute__((device)) + //===----------------------------------------------------------------------===// // Test AMDGPU builtins //===----------------------------------------------------------------------===// >From 8d47a06e2a9706b08a4f54c8d9bb36b6ac5468de Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Fri, 10 Jul 2026 20:51:23 -0400 Subject: [PATCH 6/6] remove unneccessary headers from tests --- clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip | 2 -- clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip | 2 -- 2 files changed, 4 deletions(-) diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip index aa10cd5b299b7..0ff6a5f7b83d1 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip @@ -1,5 +1,3 @@ -#include "../CodeGenCUDA/Inputs/cuda.h" - // REQUIRES: amdgpu-registered-target // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ // RUN: -target-cpu gfx1010 -fcuda-is-device -emit-cir %s -o %t.cir diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip index 90fdc2f0d9ae1..a6fb1c11e47ab 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip @@ -1,5 +1,3 @@ -#include "../CodeGenCUDA/Inputs/cuda.h" - // REQUIRES: amdgpu-registered-target // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ // RUN: -target-cpu gfx1100 -fcuda-is-device -emit-cir %s -o %t.cir _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
