[clang] [CIR][AMDGPU] Add support for AMDGCN permlane builtins (PR #197526)
Ayokunle Amodu via cfe-commits
cfe-commits at lists.llvm.org
Fri Jul 17 09:10:34 PDT 2026
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 <ayokunle321 at gmail.com>
Date: Wed, 13 May 2026 14:23:00 -0400
Subject: [PATCH 1/7] 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 <ayokunle321 at gmail.com>
Date: Wed, 13 May 2026 14:39:35 -0400
Subject: [PATCH 2/7] 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 <ayokunle321 at gmail.com>
Date: Thu, 14 May 2026 16:13:33 -0400
Subject: [PATCH 3/7] 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 <ayokunle321 at gmail.com>
Date: Fri, 10 Jul 2026 15:28:36 -0400
Subject: [PATCH 4/7] 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 <ayokunle321 at gmail.com>
Date: Fri, 10 Jul 2026 15:30:34 -0400
Subject: [PATCH 5/7] 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 <ayokunle321 at gmail.com>
Date: Fri, 10 Jul 2026 20:51:23 -0400
Subject: [PATCH 6/7] 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
>From 795864488711226a7a8ebeb83d317704a0836489 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Fri, 17 Jul 2026 12:10:21 -0400
Subject: [PATCH 7/7] Add -fclangir flag to RUN command
---
clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip
index 0ff6a5f7b83d1..83f33ef918c87 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip
@@ -19,7 +19,7 @@
// 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: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
// RUN: -target-cpu gfx1012 -fcuda-is-device -emit-llvm %s -o %t.ll
// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
More information about the cfe-commits
mailing list