[clang] [CIR][OpenCL] Lower kernel argument metadata to LLVM dialect metadata (PR #217835)
Akimasa Watanuki via cfe-commits
cfe-commits at lists.llvm.org
Wed Aug 26 07:20:40 PDT 2026
https://github.com/Men-cotton updated https://github.com/llvm/llvm-project/pull/217835
>From 881dab9f59bb3dfaf5c362784cfbd76856e684c3 Mon Sep 17 00:00:00 2001
From: mencotton <mencotton0410 at gmail.com>
Date: Mon, 8 Jun 2026 20:55:08 +0900
Subject: [PATCH 1/3] [CIR][OpenCL] Lower kernel argument metadata to LLVM
dialect metadata
Lower CIR OpenCL kernel-argument metadata into LLVMFuncOp function_metadata instead of mutating LLVM IR directly. Keep CIR storage in source-level language address spaces while emitting the fixed SPIR 2.0 metadata numbers.
---
clang/include/clang/CIR/Dialect/IR/CIROps.td | 2 +-
.../CIR/Lowering/DirectToLLVM/CMakeLists.txt | 1 +
.../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 50 +++--
.../DirectToLLVM/OpenCLMetadataLowering.cpp | 175 ++++++++++++++++++
.../DirectToLLVM/OpenCLMetadataLowering.h | 38 ++++
.../kernel-arg-info-single-as.cl | 13 ++
.../test/CIR/CodeGenOpenCL/kernel-arg-info.cl | 68 +++++++
.../CIR/CodeGenOpenCL/kernel-arg-metadata.cl | 13 ++
.../Lowering/opencl-kernel-arg-metadata.cir | 22 +++
9 files changed, 361 insertions(+), 21 deletions(-)
create mode 100644 clang/lib/CIR/Lowering/DirectToLLVM/OpenCLMetadataLowering.cpp
create mode 100644 clang/lib/CIR/Lowering/DirectToLLVM/OpenCLMetadataLowering.h
create mode 100644 clang/test/CIR/Lowering/opencl-kernel-arg-metadata.cir
diff --git a/clang/include/clang/CIR/Dialect/IR/CIROps.td b/clang/include/clang/CIR/Dialect/IR/CIROps.td
index a52c37e4860ce..79e080d268f5b 100644
--- a/clang/include/clang/CIR/Dialect/IR/CIROps.td
+++ b/clang/include/clang/CIR/Dialect/IR/CIROps.td
@@ -4354,7 +4354,7 @@ def CIR_FuncOp : CIR_Op<"func", [
static mlir::StringRef getLinkageAttrNameString() { return "linkage"; }
void lowerFuncAttributes(
- cir::FuncOp func, bool filterArgAndResAttrs,
+ cir::FuncOp func, bool includeFunctionOnlyAttrs,
mlir::SmallVectorImpl<mlir::NamedAttribute> &result) const;
mlir::LogicalResult
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt b/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt
index 3cdfdfe768c9e..3982b8db2d269 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt
@@ -8,6 +8,7 @@ get_property(dialect_libs GLOBAL PROPERTY MLIR_DIALECT_LIBS)
add_clang_library(clangCIRLoweringDirectToLLVM
LowerToLLVM.cpp
LowerToLLVMIR.cpp
+ OpenCLMetadataLowering.cpp
DEPENDS
CIRLowering
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
index 20db99f11fbef..21b69a7ec7327 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
@@ -11,6 +11,7 @@
//===----------------------------------------------------------------------===//
#include "LowerToLLVM.h"
+#include "OpenCLMetadataLowering.h"
#include <array>
#include <optional>
@@ -2677,27 +2678,33 @@ mlir::LogicalResult CIRToLLVMAbsOpLowering::matchAndRewrite(
return mlir::success();
}
-/// Convert the `cir.func` attributes to `llvm.func` attributes.
-/// Only retain those attributes that are not constructed by
-/// `LLVMFuncOp::build`. If `filterArgAttrs` is set, also filter out
-/// argument attributes.
+/// Return true for attributes constructed by `LLVMFuncOp::build`.
+static bool shouldDropFuncAttribute(cir::FuncOp func, mlir::NamedAttribute attr,
+ mlir::StringRef linkageAttrName) {
+ return attr.getName() == mlir::SymbolTable::getSymbolAttrName() ||
+ attr.getName() == func.getFunctionTypeAttrName() ||
+ attr.getName() == linkageAttrName ||
+ attr.getName() == func.getCallingConvAttrName() ||
+ attr.getName() == func.getDsoLocalAttrName() ||
+ attr.getName() == func.getInlineKindAttrName() ||
+ attr.getName() == func.getSideEffectAttrName() ||
+ attr.getName() == CIRDialect::getNoReturnAttrName() ||
+ attr.getName() == CIRDialect::getStrictFPAttrName() ||
+ attr.getName() == func.getAnnotationsAttrName();
+}
+
+/// Lower `cir.func` attributes for an `LLVMFuncOp` or `LLVM::AliasOp`.
+/// Drop attributes populated by the destination op builder. If
+/// `includeFunctionOnlyAttrs` is false, also omit attributes that are only
+/// valid on functions.
void CIRToLLVMFuncOpLowering::lowerFuncAttributes(
- cir::FuncOp func, bool filterArgAndResAttrs,
+ cir::FuncOp func, bool includeFunctionOnlyAttrs,
SmallVectorImpl<mlir::NamedAttribute> &result) const {
+ OpenCLFunctionMetadataLowering openCLMetadataLowering(func.getContext());
for (mlir::NamedAttribute attr : func->getAttrs()) {
- if (attr.getName() == mlir::SymbolTable::getSymbolAttrName() ||
- attr.getName() == func.getFunctionTypeAttrName() ||
- attr.getName() == getLinkageAttrNameString() ||
- attr.getName() == func.getCallingConvAttrName() ||
- attr.getName() == func.getDsoLocalAttrName() ||
- attr.getName() == func.getInlineKindAttrName() ||
- attr.getName() == func.getSideEffectAttrName() ||
- attr.getName() == CIRDialect::getNoReturnAttrName() ||
- attr.getName() == CIRDialect::getStrictFPAttrName() ||
- attr.getName() == func.getAnnotationsAttrName() ||
- (filterArgAndResAttrs &&
- (attr.getName() == func.getArgAttrsAttrName() ||
- attr.getName() == func.getResAttrsAttrName())))
+ if (shouldDropFuncAttribute(func, attr, getLinkageAttrNameString()))
+ continue;
+ if (openCLMetadataLowering.lower(attr, includeFunctionOnlyAttrs))
continue;
assert(!cir::MissingFeatures::opFuncExtraAttrs());
@@ -2710,13 +2717,16 @@ void CIRToLLVMFuncOpLowering::lowerFuncAttributes(
}
result.push_back(attr);
}
+
+ if (includeFunctionOnlyAttrs)
+ openCLMetadataLowering.appendAttrs(result);
}
mlir::LogicalResult CIRToLLVMFuncOpLowering::matchAndRewriteAlias(
cir::FuncOp op, llvm::StringRef aliasee, mlir::Type ty, OpAdaptor adaptor,
mlir::ConversionPatternRewriter &rewriter) const {
SmallVector<mlir::NamedAttribute, 4> attributes;
- lowerFuncAttributes(op, /*filterArgAndResAttrs=*/false, attributes);
+ lowerFuncAttributes(op, /*includeFunctionOnlyAttrs=*/false, attributes);
mlir::Location loc = op.getLoc();
auto aliasOp = rewriter.replaceOpWithNewOp<mlir::LLVM::AliasOp>(
@@ -2777,7 +2787,7 @@ mlir::LogicalResult CIRToLLVMFuncOpLowering::matchAndRewrite(
mlir::LLVM::Linkage linkage = convertLinkage(op.getLinkage());
mlir::LLVM::CConv cconv = convertCallingConv(op.getCallingConv());
SmallVector<mlir::NamedAttribute, 4> attributes;
- lowerFuncAttributes(op, /*filterArgAndResAttrs=*/false, attributes);
+ lowerFuncAttributes(op, /*includeFunctionOnlyAttrs=*/true, attributes);
mlir::LLVM::LLVMFuncOp fn = mlir::LLVM::LLVMFuncOp::create(
rewriter, loc, op.getName(), llvmFnTy, linkage, isDsoLocal, cconv,
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/OpenCLMetadataLowering.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/OpenCLMetadataLowering.cpp
new file mode 100644
index 0000000000000..fb4897b93d44f
--- /dev/null
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/OpenCLMetadataLowering.cpp
@@ -0,0 +1,175 @@
+//===- OpenCLMetadataLowering.cpp - OpenCL metadata lowering --------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "OpenCLMetadataLowering.h"
+
+#include "mlir/Dialect/LLVMIR/LLVMAttrs.h"
+#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+#include "mlir/IR/Builders.h"
+#include "clang/CIR/Dialect/IR/CIRDialect.h"
+#include "llvm/ADT/TypeSwitch.h"
+#include "llvm/Support/ErrorHandling.h"
+
+namespace cir {
+namespace direct {
+
+namespace {
+class LLVMMetadataNodeBuilder {
+public:
+ explicit LLVMMetadataNodeBuilder(mlir::MLIRContext *ctx) : ctx(ctx) {}
+
+ mlir::LLVM::MDConstantAttr getI32(unsigned value) const {
+ mlir::IntegerType intTy = mlir::IntegerType::get(ctx, 32);
+ mlir::IntegerAttr intAttr = mlir::IntegerAttr::get(intTy, value);
+ return mlir::LLVM::MDConstantAttr::get(ctx, intAttr);
+ }
+
+ mlir::LLVM::MDStringAttr getString(llvm::StringRef value) const {
+ return mlir::LLVM::MDStringAttr::get(ctx,
+ mlir::StringAttr::get(ctx, value));
+ }
+
+ mlir::LLVM::MDNodeAttr
+ getNode(llvm::ArrayRef<mlir::Attribute> metadata) const {
+ return mlir::LLVM::MDNodeAttr::get(ctx, metadata);
+ }
+
+ mlir::LLVM::MDNodeAttr getI32Node(llvm::ArrayRef<unsigned> values) const {
+ llvm::SmallVector<mlir::Attribute> metadata;
+ for (unsigned value : values)
+ metadata.push_back(getI32(value));
+ return getNode(metadata);
+ }
+
+ mlir::LLVM::MDNodeAttr getStringNode(mlir::ArrayAttr attrs) const {
+ llvm::SmallVector<mlir::Attribute> metadata;
+ for (mlir::StringAttr attr : attrs.getAsRange<mlir::StringAttr>())
+ metadata.push_back(getString(attr.getValue()));
+ return getNode(metadata);
+ }
+
+private:
+ mlir::MLIRContext *ctx;
+};
+
+using KernelArgStringMetadataGetter =
+ mlir::ArrayAttr (cir::OpenCLKernelArgMetadataAttr::*)() const;
+
+struct KernelArgStringMetadataMapping {
+ llvm::StringLiteral metadataName;
+ KernelArgStringMetadataGetter getMetadata;
+ bool optional;
+};
+
+static unsigned getOpenCLArgInfoAddressSpace(cir::LangAddressSpace as) {
+ switch (as) {
+ case cir::LangAddressSpace::Default:
+ case cir::LangAddressSpace::OffloadPrivate:
+ return 0;
+ case cir::LangAddressSpace::OffloadGlobal:
+ return 1;
+ case cir::LangAddressSpace::OffloadConstant:
+ return 2;
+ case cir::LangAddressSpace::OffloadLocal:
+ return 3;
+ case cir::LangAddressSpace::OffloadGeneric:
+ return 4;
+ case cir::LangAddressSpace::OffloadGlobalDevice:
+ return 5;
+ case cir::LangAddressSpace::OffloadGlobalHost:
+ return 6;
+ }
+ llvm_unreachable("unknown CIR language address space");
+}
+
+static mlir::LLVM::MDNodeAttr
+getAddrSpaceMetadataNode(cir::OpenCLKernelArgMetadataAttr clArgMetadata,
+ const LLVMMetadataNodeBuilder &metadataBuilder) {
+ llvm::SmallVector<unsigned> addrSpaces;
+ for (cir::LangAddressSpaceAttr addressSpace :
+ clArgMetadata.getAddrSpace().getAsRange<cir::LangAddressSpaceAttr>())
+ addrSpaces.push_back(getOpenCLArgInfoAddressSpace(addressSpace.getValue()));
+ return metadataBuilder.getI32Node(addrSpaces);
+}
+
+static void addOpenCLKernelArgFunctionMetadata(
+ mlir::MLIRContext *ctx, llvm::SmallVectorImpl<mlir::Attribute> &entries,
+ llvm::StringRef name, mlir::LLVM::MDNodeAttr node) {
+ entries.push_back(mlir::LLVM::FunctionMetadataAttr::get(
+ ctx, mlir::StringAttr::get(ctx, name), node));
+}
+
+} // namespace
+
+static void convertOpenCLKernelArgMetadata(
+ cir::OpenCLKernelArgMetadataAttr clArgMetadata,
+ llvm::SmallVectorImpl<mlir::Attribute> &entries) {
+ mlir::MLIRContext *ctx = clArgMetadata.getContext();
+ LLVMMetadataNodeBuilder metadataBuilder(ctx);
+
+ addOpenCLKernelArgFunctionMetadata(
+ ctx, entries, "kernel_arg_addr_space",
+ getAddrSpaceMetadataNode(clArgMetadata, metadataBuilder));
+
+ static constexpr KernelArgStringMetadataMapping stringMetadataMappings[] = {
+ {"kernel_arg_access_qual",
+ &cir::OpenCLKernelArgMetadataAttr::getAccessQual,
+ /*optional=*/false},
+ {"kernel_arg_type", &cir::OpenCLKernelArgMetadataAttr::getType,
+ /*optional=*/false},
+ {"kernel_arg_base_type", &cir::OpenCLKernelArgMetadataAttr::getBaseType,
+ /*optional=*/false},
+ {"kernel_arg_type_qual", &cir::OpenCLKernelArgMetadataAttr::getTypeQual,
+ /*optional=*/false},
+ {"kernel_arg_name", &cir::OpenCLKernelArgMetadataAttr::getName,
+ /*optional=*/true},
+ };
+
+ for (const KernelArgStringMetadataMapping &mapping : stringMetadataMappings) {
+ mlir::ArrayAttr metadata = (clArgMetadata.*mapping.getMetadata)();
+ if (mapping.optional && !metadata)
+ continue;
+ addOpenCLKernelArgFunctionMetadata(ctx, entries, mapping.metadataName,
+ metadataBuilder.getStringNode(metadata));
+ }
+}
+
+OpenCLFunctionMetadataLowering::OpenCLFunctionMetadataLowering(
+ mlir::MLIRContext *ctx)
+ : ctx(ctx) {}
+
+bool OpenCLFunctionMetadataLowering::lower(mlir::NamedAttribute attr,
+ bool includeFunctionOnlyAttrs) {
+ return llvm::TypeSwitch<mlir::Attribute, bool>(attr.getValue())
+ .Case<cir::OpenCLKernelArgMetadataAttr>(
+ [&](cir::OpenCLKernelArgMetadataAttr clArgMetadata) {
+ if (!includeFunctionOnlyAttrs)
+ return true;
+ return lower(clArgMetadata);
+ })
+ .Default(false);
+}
+
+void OpenCLFunctionMetadataLowering::appendAttrs(
+ llvm::SmallVectorImpl<mlir::NamedAttribute> &result) const {
+ if (!functionMetadata.empty()) {
+ result.push_back(mlir::NamedAttribute(
+ mlir::LLVM::LLVMFuncOp::getFunctionMetadataAttrName(mlir::OperationName(
+ mlir::LLVM::LLVMFuncOp::getOperationName(), ctx)),
+ mlir::ArrayAttr::get(ctx, functionMetadata)));
+ }
+}
+
+bool OpenCLFunctionMetadataLowering::lower(
+ cir::OpenCLKernelArgMetadataAttr clArgMetadata) {
+ convertOpenCLKernelArgMetadata(clArgMetadata, functionMetadata);
+ return true;
+}
+
+} // namespace direct
+} // namespace cir
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/OpenCLMetadataLowering.h b/clang/lib/CIR/Lowering/DirectToLLVM/OpenCLMetadataLowering.h
new file mode 100644
index 0000000000000..d45123cae4f3f
--- /dev/null
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/OpenCLMetadataLowering.h
@@ -0,0 +1,38 @@
+//===- OpenCLMetadataLowering.h - OpenCL metadata lowering ------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef CLANG_CIR_LOWERING_DIRECTTOLLVM_OPENCLMETADATALOWERING_H
+#define CLANG_CIR_LOWERING_DIRECTTOLLVM_OPENCLMETADATALOWERING_H
+
+#include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/MLIRContext.h"
+#include "mlir/IR/OperationSupport.h"
+#include "clang/CIR/Dialect/IR/CIRAttrs.h"
+#include "llvm/ADT/SmallVector.h"
+
+namespace cir {
+namespace direct {
+
+class OpenCLFunctionMetadataLowering {
+public:
+ OpenCLFunctionMetadataLowering(mlir::MLIRContext *ctx);
+
+ bool lower(mlir::NamedAttribute attr, bool includeFunctionOnlyAttrs);
+ void appendAttrs(llvm::SmallVectorImpl<mlir::NamedAttribute> &result) const;
+
+private:
+ bool lower(cir::OpenCLKernelArgMetadataAttr clArgMetadata);
+
+ mlir::MLIRContext *ctx;
+ llvm::SmallVector<mlir::Attribute> functionMetadata;
+};
+
+} // namespace direct
+} // namespace cir
+
+#endif // CLANG_CIR_LOWERING_DIRECTTOLLVM_OPENCLMETADATALOWERING_H
diff --git a/clang/test/CIR/CodeGenOpenCL/kernel-arg-info-single-as.cl b/clang/test/CIR/CodeGenOpenCL/kernel-arg-info-single-as.cl
index 4f17d39ccb974..c60c36dc61c89 100644
--- a/clang/test/CIR/CodeGenOpenCL/kernel-arg-info-single-as.cl
+++ b/clang/test/CIR/CodeGenOpenCL/kernel-arg-info-single-as.cl
@@ -2,6 +2,10 @@
// spaces even if the target has only one address space like x86_64 does.
// RUN: %clang_cc1 %s -fclangir -cl-std=CL2.0 -triple x86_64-unknown-linux-gnu -emit-cir -o %t.cir
// RUN: FileCheck %s --input-file=%t.cir --check-prefix=CIR
+// RUN: %clang_cc1 %s -fclangir -cl-std=CL2.0 -triple x86_64-unknown-linux-gnu -emit-llvm -o %t.ll
+// RUN: FileCheck %s --input-file=%t.ll --check-prefix=LLVM
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -triple x86_64-unknown-linux-gnu -emit-llvm -o %t.ogcg.ll
+// RUN: FileCheck %s --input-file=%t.ogcg.ll --check-prefix=LLVM
kernel void spir_addr_space_kernel_args(__global int *G, __constant int *C,
__local int *L) {
@@ -11,6 +15,9 @@ kernel void spir_addr_space_kernel_args(__global int *G, __constant int *C,
// CIR-LABEL: cir.func{{.*}} @spir_addr_space_kernel_args
// CIR-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata<addr_space = [#cir<lang_address_space(offload_global)>, #cir<lang_address_space(offload_constant)>, #cir<lang_address_space(offload_local)>]
+// LLVM-DAG: define{{.*}} void @spir_addr_space_kernel_args{{.*}} !kernel_arg_addr_space ![[ADDR_SPACES:[0-9]+]]
+// LLVM-DAG: ![[ADDR_SPACES]] = !{i32 1, i32 2, i32 3}
+
kernel void global_device_host_kernel_args(
__attribute__((opencl_global_device)) int *D,
__attribute__((opencl_global_host)) int *H) {}
@@ -18,6 +25,9 @@ kernel void global_device_host_kernel_args(
// CIR-LABEL: cir.func{{.*}} @global_device_host_kernel_args
// CIR-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata<addr_space = [#cir<lang_address_space(offload_global_device)>, #cir<lang_address_space(offload_global_host)>]
+// LLVM-DAG: define{{.*}} void @global_device_host_kernel_args{{.*}} !kernel_arg_addr_space ![[GLOBAL_DEVICE_HOST_ADDR_SPACES:[0-9]+]]
+// LLVM-DAG: ![[GLOBAL_DEVICE_HOST_ADDR_SPACES]] = !{i32 5, i32 6}
+
// Target-specific address spaces stay on pointer types but do not represent an
// OpenCL language address-space qualifier.
kernel void target_address_space_kernel_arg(
@@ -26,3 +36,6 @@ kernel void target_address_space_kernel_arg(
// CIR-LABEL: cir.func{{.*}} @target_address_space_kernel_arg
// CIR-SAME: !cir.ptr<!s32i, target_address_space(5)>
// CIR-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata<addr_space = [#cir<lang_address_space(default)>]
+
+// LLVM-DAG: define{{.*}} void @target_address_space_kernel_arg{{.*}}ptr addrspace(5){{.*}} !kernel_arg_addr_space ![[TARGET_ADDR_SPACE:[0-9]+]]
+// LLVM-DAG: ![[TARGET_ADDR_SPACE]] = !{i32 0}
diff --git a/clang/test/CIR/CodeGenOpenCL/kernel-arg-info.cl b/clang/test/CIR/CodeGenOpenCL/kernel-arg-info.cl
index 7788195157715..a00bfec52b122 100644
--- a/clang/test/CIR/CodeGenOpenCL/kernel-arg-info.cl
+++ b/clang/test/CIR/CodeGenOpenCL/kernel-arg-info.cl
@@ -4,6 +4,19 @@
// RUN: %clang_cc1 %s -fclangir -cl-std=CL2.0 -triple spirv64-unknown-unknown -emit-cir -cl-kernel-arg-info -o %t.arginfo.cir
// RUN: FileCheck %s --input-file=%t.arginfo.cir --check-prefix=CIR-ARGINFO
+// RUN: %clang_cc1 %s -fclangir -cl-std=CL2.0 -triple spirv64-unknown-unknown -emit-llvm -o %t.ll
+// RUN: FileCheck %s --input-file=%t.ll --check-prefix=LLVM
+// RUN: FileCheck %s --input-file=%t.ll --check-prefix=LLVM-NO-ARGINFO
+// RUN: %clang_cc1 %s -fclangir -cl-std=CL2.0 -triple spirv64-unknown-unknown -emit-llvm -cl-kernel-arg-info -o %t.arginfo.ll
+// RUN: FileCheck %s --input-file=%t.arginfo.ll --check-prefix=LLVM
+// RUN: FileCheck %s --input-file=%t.arginfo.ll --check-prefix=LLVM-ARGINFO
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -triple spirv64-unknown-unknown -emit-llvm -o %t.ogcg.ll
+// RUN: FileCheck %s --input-file=%t.ogcg.ll --check-prefix=LLVM
+// RUN: FileCheck %s --input-file=%t.ogcg.ll --check-prefix=LLVM-NO-ARGINFO
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -triple spirv64-unknown-unknown -emit-llvm -cl-kernel-arg-info -o %t.ogcg.arginfo.ll
+// RUN: FileCheck %s --input-file=%t.ogcg.arginfo.ll --check-prefix=LLVM
+// RUN: FileCheck %s --input-file=%t.ogcg.arginfo.ll --check-prefix=LLVM-ARGINFO
+
kernel void global_qualifier_kernel_args(
global int *globalintp, global int *restrict globalintrestrictp,
global const int *globalconstintp,
@@ -29,6 +42,16 @@ kernel void global_qualifier_kernel_args(
// CIR-ARGINFO-SAME: type_qual = ["", "restrict", "const", "restrict const", "const volatile", "restrict const volatile", "volatile", "restrict volatile"]
// CIR-ARGINFO-SAME: name = ["globalintp", "globalintrestrictp", "globalconstintp", "globalconstintrestrictp", "globalconstvolatileintp", "globalconstvolatileintrestrictp", "globalvolatileintp", "globalvolatileintrestrictp"]
+// LLVM-DAG: define{{.*}} void @global_qualifier_kernel_args{{.*}} !kernel_arg_addr_space ![[GLOBAL_ADDR_SPACES:[0-9]+]] !kernel_arg_access_qual ![[GLOBAL_ACCESS_QUALS:[0-9]+]] !kernel_arg_type ![[GLOBAL_ARG_TYPES:[0-9]+]] !kernel_arg_base_type ![[GLOBAL_ARG_TYPES]] !kernel_arg_type_qual ![[GLOBAL_TYPE_QUALS:[0-9]+]]
+// LLVM-DAG: ![[GLOBAL_ADDR_SPACES]] = !{i32 1, i32 1, i32 1, i32 1, i32 1, i32 1, i32 1, i32 1}
+// LLVM-DAG: ![[GLOBAL_ACCESS_QUALS]] = !{!"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none"}
+// LLVM-DAG: ![[GLOBAL_ARG_TYPES]] = !{!"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*"}
+// LLVM-DAG: ![[GLOBAL_TYPE_QUALS]] = !{!"", !"restrict", !"const", !"restrict const", !"const volatile", !"restrict const volatile", !"volatile", !"restrict volatile"}
+// LLVM-ARGINFO-DAG: define{{.*}} void @global_qualifier_kernel_args{{.*}} !kernel_arg_name ![[GLOBAL_ARG_NAMES:[0-9]+]]
+// LLVM-ARGINFO-DAG: ![[GLOBAL_ARG_NAMES]] = !{!"globalintp", !"globalintrestrictp", !"globalconstintp", !"globalconstintrestrictp", !"globalconstvolatileintp", !"globalconstvolatileintrestrictp", !"globalvolatileintp", !"globalvolatileintrestrictp"}
+// LLVM-NO-ARGINFO: define{{.*}} void @global_qualifier_kernel_args
+// LLVM-NO-ARGINFO-NOT: !kernel_arg_name
+
kernel void constant_kernel_args(constant int *constantintp,
constant int *restrict constantintrestrictp) {}
@@ -48,6 +71,14 @@ kernel void constant_kernel_args(constant int *constantintp,
// CIR-ARGINFO-SAME: type_qual = ["const", "restrict const"]
// CIR-ARGINFO-SAME: name = ["constantintp", "constantintrestrictp"]
+// LLVM-DAG: define{{.*}} void @constant_kernel_args{{.*}} !kernel_arg_addr_space ![[CONSTANT_ADDR_SPACES:[0-9]+]] !kernel_arg_access_qual ![[CONSTANT_ACCESS_QUALS:[0-9]+]] !kernel_arg_type ![[CONSTANT_ARG_TYPES:[0-9]+]] !kernel_arg_base_type ![[CONSTANT_ARG_TYPES]] !kernel_arg_type_qual ![[CONSTANT_TYPE_QUALS:[0-9]+]]
+// LLVM-DAG: ![[CONSTANT_ADDR_SPACES]] = !{i32 2, i32 2}
+// LLVM-DAG: ![[CONSTANT_ACCESS_QUALS]] = !{!"none", !"none"}
+// LLVM-DAG: ![[CONSTANT_ARG_TYPES]] = !{!"int*", !"int*"}
+// LLVM-DAG: ![[CONSTANT_TYPE_QUALS]] = !{!"const", !"restrict const"}
+// LLVM-ARGINFO-DAG: define{{.*}} void @constant_kernel_args{{.*}} !kernel_arg_name ![[CONSTANT_ARG_NAMES:[0-9]+]]
+// LLVM-ARGINFO-DAG: ![[CONSTANT_ARG_NAMES]] = !{!"constantintp", !"constantintrestrictp"}
+
kernel void local_qualifier_kernel_args(
local int *localintp, local int *restrict localintrestrictp,
local const int *localconstintp,
@@ -73,6 +104,11 @@ kernel void local_qualifier_kernel_args(
// CIR-ARGINFO-SAME: type_qual = ["", "restrict", "const", "restrict const", "const volatile", "restrict const volatile", "volatile", "restrict volatile"]
// CIR-ARGINFO-SAME: name = ["localintp", "localintrestrictp", "localconstintp", "localconstintrestrictp", "localconstvolatileintp", "localconstvolatileintrestrictp", "localvolatileintp", "localvolatileintrestrictp"]
+// LLVM-DAG: define{{.*}} void @local_qualifier_kernel_args{{.*}} !kernel_arg_addr_space ![[LOCAL_ADDR_SPACES:[0-9]+]] !kernel_arg_access_qual ![[GLOBAL_ACCESS_QUALS]] !kernel_arg_type ![[GLOBAL_ARG_TYPES]] !kernel_arg_base_type ![[GLOBAL_ARG_TYPES]] !kernel_arg_type_qual ![[GLOBAL_TYPE_QUALS]]
+// LLVM-DAG: ![[LOCAL_ADDR_SPACES]] = !{i32 3, i32 3, i32 3, i32 3, i32 3, i32 3, i32 3, i32 3}
+// LLVM-ARGINFO-DAG: define{{.*}} void @local_qualifier_kernel_args{{.*}} !kernel_arg_name ![[LOCAL_ARG_NAMES:[0-9]+]]
+// LLVM-ARGINFO-DAG: ![[LOCAL_ARG_NAMES]] = !{!"localintp", !"localintrestrictp", !"localconstintp", !"localconstintrestrictp", !"localconstvolatileintp", !"localconstvolatileintrestrictp", !"localvolatileintp", !"localvolatileintrestrictp"}
+
kernel void private_qualifier_kernel_args(int X, const int constint,
const volatile int constvolatileint,
volatile int volatileint) {}
@@ -93,6 +129,14 @@ kernel void private_qualifier_kernel_args(int X, const int constint,
// CIR-ARGINFO-SAME: type_qual = ["", "", "", ""]
// CIR-ARGINFO-SAME: name = ["X", "constint", "constvolatileint", "volatileint"]
+// LLVM-DAG: define{{.*}} void @private_qualifier_kernel_args{{.*}} !kernel_arg_addr_space ![[PRIVATE_ADDR_SPACES:[0-9]+]] !kernel_arg_access_qual ![[PRIVATE_ACCESS_QUALS:[0-9]+]] !kernel_arg_type ![[PRIVATE_ARG_TYPES:[0-9]+]] !kernel_arg_base_type ![[PRIVATE_ARG_TYPES]] !kernel_arg_type_qual ![[PRIVATE_TYPE_QUALS:[0-9]+]]
+// LLVM-DAG: ![[PRIVATE_ADDR_SPACES]] = !{i32 0, i32 0, i32 0, i32 0}
+// LLVM-DAG: ![[PRIVATE_ACCESS_QUALS]] = !{!"none", !"none", !"none", !"none"}
+// LLVM-DAG: ![[PRIVATE_ARG_TYPES]] = !{!"int", !"int", !"int", !"int"}
+// LLVM-DAG: ![[PRIVATE_TYPE_QUALS]] = !{!"", !"", !"", !""}
+// LLVM-ARGINFO-DAG: define{{.*}} void @private_qualifier_kernel_args{{.*}} !kernel_arg_name ![[PRIVATE_ARG_NAMES:[0-9]+]]
+// LLVM-ARGINFO-DAG: ![[PRIVATE_ARG_NAMES]] = !{!"X", !"constint", !"constvolatileint", !"volatileint"}
+
typedef unsigned int myunsignedint;
kernel void typedef_kernel_args(__global unsigned int *X,
__global myunsignedint *Y) {}
@@ -113,6 +157,14 @@ kernel void typedef_kernel_args(__global unsigned int *X,
// CIR-ARGINFO-SAME: type_qual = ["", ""]
// CIR-ARGINFO-SAME: name = ["X", "Y"]
+// LLVM-DAG: define{{.*}} void @typedef_kernel_args{{.*}} !kernel_arg_addr_space ![[TYPEDEF_ADDR_SPACES:[0-9]+]] !kernel_arg_access_qual ![[CONSTANT_ACCESS_QUALS]] !kernel_arg_type ![[TYPEDEF_ARG_TYPES:[0-9]+]] !kernel_arg_base_type ![[TYPEDEF_BASE_TYPES:[0-9]+]] !kernel_arg_type_qual ![[TYPEDEF_TYPE_QUALS:[0-9]+]]
+// LLVM-DAG: ![[TYPEDEF_ADDR_SPACES]] = !{i32 1, i32 1}
+// LLVM-DAG: ![[TYPEDEF_ARG_TYPES]] = !{!"uint*", !"myunsignedint*"}
+// LLVM-DAG: ![[TYPEDEF_BASE_TYPES]] = !{!"uint*", !"uint*"}
+// LLVM-DAG: ![[TYPEDEF_TYPE_QUALS]] = !{!"", !""}
+// LLVM-ARGINFO-DAG: define{{.*}} void @typedef_kernel_args{{.*}} !kernel_arg_name ![[TYPEDEF_ARG_NAMES:[0-9]+]]
+// LLVM-ARGINFO-DAG: ![[TYPEDEF_ARG_NAMES]] = !{!"X", !"Y"}
+
typedef char char16 __attribute__((ext_vector_type(16)));
__kernel void vector_typedef_kernel_arg(__global char16 arg[]) {}
@@ -132,6 +184,15 @@ __kernel void vector_typedef_kernel_arg(__global char16 arg[]) {}
// CIR-ARGINFO-SAME: type_qual = [""]
// CIR-ARGINFO-SAME: name = ["arg"]
+// LLVM-DAG: define{{.*}} void @vector_typedef_kernel_arg{{.*}} !kernel_arg_addr_space ![[VECTOR_TYPEDEF_ADDR_SPACES:[0-9]+]] !kernel_arg_access_qual ![[VECTOR_TYPEDEF_ACCESS_QUALS:[0-9]+]] !kernel_arg_type ![[VECTOR_TYPEDEF_ARG_TYPES:[0-9]+]] !kernel_arg_base_type ![[VECTOR_TYPEDEF_BASE_TYPES:[0-9]+]] !kernel_arg_type_qual ![[VECTOR_TYPEDEF_TYPE_QUALS:[0-9]+]]
+// LLVM-DAG: ![[VECTOR_TYPEDEF_ADDR_SPACES]] = !{i32 1}
+// LLVM-DAG: ![[VECTOR_TYPEDEF_ACCESS_QUALS]] = !{!"none"}
+// LLVM-DAG: ![[VECTOR_TYPEDEF_ARG_TYPES]] = !{!"char16*"}
+// LLVM-DAG: ![[VECTOR_TYPEDEF_BASE_TYPES]] = !{!"char __attribute__((ext_vector_type(16)))*"}
+// LLVM-DAG: ![[VECTOR_TYPEDEF_TYPE_QUALS]] = !{!""}
+// LLVM-ARGINFO-DAG: define{{.*}} void @vector_typedef_kernel_arg{{.*}} !kernel_arg_name ![[VECTOR_TYPEDEF_ARG_NAMES:[0-9]+]]
+// LLVM-ARGINFO-DAG: ![[VECTOR_TYPEDEF_ARG_NAMES]] = !{!"arg"}
+
kernel void signed_char_kernel_args(signed char sc1,
global const signed char *sc2) {}
@@ -150,3 +211,10 @@ kernel void signed_char_kernel_args(signed char sc1,
// CIR-ARGINFO-SAME: base_type = ["char", "char*"]
// CIR-ARGINFO-SAME: type_qual = ["", "const"]
// CIR-ARGINFO-SAME: name = ["sc1", "sc2"]
+
+// LLVM-DAG: define{{.*}} void @signed_char_kernel_args{{.*}} !kernel_arg_addr_space ![[SIGNED_CHAR_ADDR_SPACES:[0-9]+]] !kernel_arg_access_qual ![[CONSTANT_ACCESS_QUALS]] !kernel_arg_type ![[SIGNED_CHAR_ARG_TYPES:[0-9]+]] !kernel_arg_base_type ![[SIGNED_CHAR_ARG_TYPES]] !kernel_arg_type_qual ![[SIGNED_CHAR_TYPE_QUALS:[0-9]+]]
+// LLVM-DAG: ![[SIGNED_CHAR_ADDR_SPACES]] = !{i32 0, i32 1}
+// LLVM-DAG: ![[SIGNED_CHAR_ARG_TYPES]] = !{!"char", !"char*"}
+// LLVM-DAG: ![[SIGNED_CHAR_TYPE_QUALS]] = !{!"", !"const"}
+// LLVM-ARGINFO-DAG: define{{.*}} void @signed_char_kernel_args{{.*}} !kernel_arg_name ![[SIGNED_CHAR_ARG_NAMES:[0-9]+]]
+// LLVM-ARGINFO-DAG: ![[SIGNED_CHAR_ARG_NAMES]] = !{!"sc1", !"sc2"}
diff --git a/clang/test/CIR/CodeGenOpenCL/kernel-arg-metadata.cl b/clang/test/CIR/CodeGenOpenCL/kernel-arg-metadata.cl
index b1ae2d8250b69..b97b47ff78dd2 100644
--- a/clang/test/CIR/CodeGenOpenCL/kernel-arg-metadata.cl
+++ b/clang/test/CIR/CodeGenOpenCL/kernel-arg-metadata.cl
@@ -1,12 +1,25 @@
// RUN: %clang_cc1 %s -fclangir -triple spirv64-unknown-unknown -emit-cir -o %t.cir
// RUN: FileCheck %s --input-file=%t.cir --check-prefix=CIR
+// RUN: %clang_cc1 %s -fclangir -triple spirv64-unknown-unknown -emit-llvm -o %t.ll
+// RUN: FileCheck %s --input-file=%t.ll --check-prefix=LLVM
+// RUN: %clang_cc1 %s -triple spirv64-unknown-unknown -emit-llvm -o %t.ogcg.ll
+// RUN: FileCheck %s --input-file=%t.ogcg.ll --check-prefix=LLVM
extern __kernel void alias_kernel_function(void)
__attribute__((alias("kernel_function")));
// CIR-LABEL: cir.func @alias_kernel_function() alias(@kernel_function)
+// LLVM-LABEL: @alias_kernel_function = alias void (), ptr @kernel_function
+// LLVM-NOT: !kernel_arg_
__kernel void kernel_function() {}
// CIR-LABEL: cir.func @kernel_function()
// CIR-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata<addr_space = [], access_qual = [], type = [], base_type = [], type_qual = []>
+// LLVM-LABEL: define spir_kernel void @kernel_function()
+// LLVM-SAME: !kernel_arg_addr_space ![[EMPTY_ARG_METADATA:[0-9]+]]
+// LLVM-SAME: !kernel_arg_access_qual ![[EMPTY_ARG_METADATA]]
+// LLVM-SAME: !kernel_arg_type ![[EMPTY_ARG_METADATA]]
+// LLVM-SAME: !kernel_arg_base_type ![[EMPTY_ARG_METADATA]]
+// LLVM-SAME: !kernel_arg_type_qual ![[EMPTY_ARG_METADATA]]
+// LLVM: ![[EMPTY_ARG_METADATA]] = !{}
diff --git a/clang/test/CIR/Lowering/opencl-kernel-arg-metadata.cir b/clang/test/CIR/Lowering/opencl-kernel-arg-metadata.cir
new file mode 100644
index 0000000000000..11541fde880cf
--- /dev/null
+++ b/clang/test/CIR/Lowering/opencl-kernel-arg-metadata.cir
@@ -0,0 +1,22 @@
+// RUN: cir-opt %s -cir-to-llvm -o - | FileCheck %s --check-prefix=MLIR
+
+module attributes {cir.triple = "x86_64-unknown-linux-gnu"} {
+ cir.func @opencl_kernel_with_names() attributes {cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata<addr_space = [#cir<lang_address_space(offload_global)>, #cir<lang_address_space(default)>], access_qual = ["none", "none"], type = ["uint*", "int"], base_type = ["uint*", "int"], type_qual = ["restrict", ""], name = ["data", "count"]>} {
+ cir.return
+ }
+ cir.func @opencl_target() {
+ cir.return
+ }
+ cir.func @opencl_alias() alias(@opencl_target) attributes {cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata<addr_space = [#cir<lang_address_space(offload_global)>], access_qual = ["none"], type = ["uint*"], base_type = ["uint*"], type_qual = [""]>}
+}
+
+// MLIR-LABEL: llvm.func @opencl_kernel_with_names()
+// MLIR-SAME: function_metadata =
+// MLIR-SAME: #llvm.func_metadata<"kernel_arg_addr_space", <#llvm.md_const<1 : i32>, #llvm.md_const<0 : i32>>>
+// MLIR-SAME: #llvm.func_metadata<"kernel_arg_access_qual", <#llvm.md_string<"none">, #llvm.md_string<"none">>>
+// MLIR-SAME: #llvm.func_metadata<"kernel_arg_type", <#llvm.md_string<"uint*">, #llvm.md_string<"int">>>
+// MLIR-SAME: #llvm.func_metadata<"kernel_arg_base_type", <#llvm.md_string<"uint*">, #llvm.md_string<"int">>>
+// MLIR-SAME: #llvm.func_metadata<"kernel_arg_type_qual", <#llvm.md_string<"restrict">, #llvm.md_string<"">>>
+// MLIR-SAME: #llvm.func_metadata<"kernel_arg_name", <#llvm.md_string<"data">, #llvm.md_string<"count">>>
+// MLIR-LABEL: llvm.mlir.alias external @opencl_alias
+// MLIR-NOT: function_metadata
>From 5ce5dc85a6fba58d0a46e5ae27a3a19753167563 Mon Sep 17 00:00:00 2001
From: mencotton <mencotton0410 at gmail.com>
Date: Wed, 26 Aug 2026 12:37:38 +0900
Subject: [PATCH 2/3] fix: rename OpenCL metadata lowering files
---
clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt | 2 +-
clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 2 +-
...MetadataLowering.cpp => LowerToLLVMOpenCLMetadata.cpp} | 4 ++--
...enCLMetadataLowering.h => LowerToLLVMOpenCLMetadata.h} | 8 ++++----
4 files changed, 8 insertions(+), 8 deletions(-)
rename clang/lib/CIR/Lowering/DirectToLLVM/{OpenCLMetadataLowering.cpp => LowerToLLVMOpenCLMetadata.cpp} (98%)
rename clang/lib/CIR/Lowering/DirectToLLVM/{OpenCLMetadataLowering.h => LowerToLLVMOpenCLMetadata.h} (78%)
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt b/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt
index 3982b8db2d269..26d9f65f865be 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt
@@ -8,7 +8,7 @@ get_property(dialect_libs GLOBAL PROPERTY MLIR_DIALECT_LIBS)
add_clang_library(clangCIRLoweringDirectToLLVM
LowerToLLVM.cpp
LowerToLLVMIR.cpp
- OpenCLMetadataLowering.cpp
+ LowerToLLVMOpenCLMetadata.cpp
DEPENDS
CIRLowering
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
index 21b69a7ec7327..c3f36502c63d4 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
@@ -11,7 +11,7 @@
//===----------------------------------------------------------------------===//
#include "LowerToLLVM.h"
-#include "OpenCLMetadataLowering.h"
+#include "LowerToLLVMOpenCLMetadata.h"
#include <array>
#include <optional>
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/OpenCLMetadataLowering.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.cpp
similarity index 98%
rename from clang/lib/CIR/Lowering/DirectToLLVM/OpenCLMetadataLowering.cpp
rename to clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.cpp
index fb4897b93d44f..57fce3ec00fc3 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/OpenCLMetadataLowering.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.cpp
@@ -1,4 +1,4 @@
-//===- OpenCLMetadataLowering.cpp - OpenCL metadata lowering --------------===//
+//===- LowerToLLVMOpenCLMetadata.cpp - OpenCL metadata lowering -----------===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
@@ -6,7 +6,7 @@
//
//===----------------------------------------------------------------------===//
-#include "OpenCLMetadataLowering.h"
+#include "LowerToLLVMOpenCLMetadata.h"
#include "mlir/Dialect/LLVMIR/LLVMAttrs.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/OpenCLMetadataLowering.h b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.h
similarity index 78%
rename from clang/lib/CIR/Lowering/DirectToLLVM/OpenCLMetadataLowering.h
rename to clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.h
index d45123cae4f3f..ddc3b433663d5 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/OpenCLMetadataLowering.h
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.h
@@ -1,4 +1,4 @@
-//===- OpenCLMetadataLowering.h - OpenCL metadata lowering ------*- C++ -*-===//
+//===- LowerToLLVMOpenCLMetadata.h - OpenCL metadata lowering ---*- C++ -*-===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
@@ -6,8 +6,8 @@
//
//===----------------------------------------------------------------------===//
-#ifndef CLANG_CIR_LOWERING_DIRECTTOLLVM_OPENCLMETADATALOWERING_H
-#define CLANG_CIR_LOWERING_DIRECTTOLLVM_OPENCLMETADATALOWERING_H
+#ifndef CLANG_CIR_LOWERING_DIRECTTOLLVM_LOWERTOLLVMOPENCLMETADATA_H
+#define CLANG_CIR_LOWERING_DIRECTTOLLVM_LOWERTOLLVMOPENCLMETADATA_H
#include "mlir/IR/BuiltinOps.h"
#include "mlir/IR/MLIRContext.h"
@@ -35,4 +35,4 @@ class OpenCLFunctionMetadataLowering {
} // namespace direct
} // namespace cir
-#endif // CLANG_CIR_LOWERING_DIRECTTOLLVM_OPENCLMETADATALOWERING_H
+#endif // CLANG_CIR_LOWERING_DIRECTTOLLVM_LOWERTOLLVMOPENCLMETADATA_H
>From 5d5a58bc46fe6df4a84f4a85e46866d37d4a6c2a Mon Sep 17 00:00:00 2001
From: mencotton <mencotton0410 at gmail.com>
Date: Wed, 26 Aug 2026 12:39:18 +0900
Subject: [PATCH 3/3] fix: make typed `lower` overloads return void
---
.../CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.cpp | 6 +++---
.../CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.h | 2 +-
2 files changed, 4 insertions(+), 4 deletions(-)
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.cpp
index 57fce3ec00fc3..e2e7188df009b 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.cpp
@@ -150,7 +150,8 @@ bool OpenCLFunctionMetadataLowering::lower(mlir::NamedAttribute attr,
[&](cir::OpenCLKernelArgMetadataAttr clArgMetadata) {
if (!includeFunctionOnlyAttrs)
return true;
- return lower(clArgMetadata);
+ lower(clArgMetadata);
+ return true;
})
.Default(false);
}
@@ -165,10 +166,9 @@ void OpenCLFunctionMetadataLowering::appendAttrs(
}
}
-bool OpenCLFunctionMetadataLowering::lower(
+void OpenCLFunctionMetadataLowering::lower(
cir::OpenCLKernelArgMetadataAttr clArgMetadata) {
convertOpenCLKernelArgMetadata(clArgMetadata, functionMetadata);
- return true;
}
} // namespace direct
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.h b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.h
index ddc3b433663d5..3865711656c35 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.h
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.h
@@ -26,7 +26,7 @@ class OpenCLFunctionMetadataLowering {
void appendAttrs(llvm::SmallVectorImpl<mlir::NamedAttribute> &result) const;
private:
- bool lower(cir::OpenCLKernelArgMetadataAttr clArgMetadata);
+ void lower(cir::OpenCLKernelArgMetadataAttr clArgMetadata);
mlir::MLIRContext *ctx;
llvm::SmallVector<mlir::Attribute> functionMetadata;
More information about the cfe-commits
mailing list