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 <[email protected]> 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 <[email protected]> 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 <[email protected]> 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; _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
