[llvm-branch-commits] [clang] [CIR][OpenCL] Lower kernel argument metadata to LLVM IR (PR #200582)
Akimasa Watanuki via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Sat May 30 07:39:25 PDT 2026
https://github.com/Men-cotton created https://github.com/llvm/llvm-project/pull/200582
Translate CIR OpenCL kernel argument metadata into the LLVM IR kernel_arg_* metadata attached to kernel functions. Preserve optional argument names so -cl-kernel-arg-info controls the LLVM metadata surface through the CIR attribute.
>From c7cd58a9e6dbdd4c0a0850b411b143df312fa775 Mon Sep 17 00:00:00 2001
From: mencotton <mencotton0410 at gmail.com>
Date: Sun, 24 May 2026 00:57:42 +0900
Subject: [PATCH] [CIR][OpenCL] Lower kernel argument metadata to LLVM IR
Translate CIR OpenCL kernel argument metadata into the LLVM IR kernel_arg_* metadata attached to kernel functions. Preserve optional argument names so -cl-kernel-arg-info controls the LLVM metadata surface through the CIR attribute.
---
.../Lowering/DirectToLLVM/LowerToLLVMIR.cpp | 75 ++++++++++++++++++-
.../kernel-arg-info-single-as.cl | 12 +++
.../test/CIR/CodeGenOpenCL/kernel-arg-info.cl | 60 +++++++++++++++
.../CIR/CodeGenOpenCL/kernel-arg-metadata.cl | 12 +++
4 files changed, 158 insertions(+), 1 deletion(-)
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp
index dbcd0aed88056..fd420cf3153bf 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp
@@ -15,10 +15,14 @@
#include "mlir/IR/DialectRegistry.h"
#include "mlir/Target/LLVMIR/LLVMTranslationInterface.h"
#include "mlir/Target/LLVMIR/ModuleTranslation.h"
+#include "clang/CIR/Dialect/IR/CIRAttrs.h"
#include "clang/CIR/Dialect/IR/CIRDialect.h"
#include "llvm/ADT/ArrayRef.h"
#include "llvm/IR/Constant.h"
+#include "llvm/IR/Function.h"
#include "llvm/IR/GlobalVariable.h"
+#include "llvm/IR/Metadata.h"
+#include "llvm/Support/ErrorHandling.h"
using namespace llvm;
@@ -75,11 +79,80 @@ class CIRDialectLLVMIRTranslationInterface
// Strip the "cir." prefix to get the LLVM attribute name.
llvm::StringRef llvmAttrName = attrName.substr(strlen("cir."));
- if (auto strAttr = mlir::dyn_cast<mlir::StringAttr>(attribute.getValue()))
+ if (auto clArgMetadata = mlir::dyn_cast<cir::OpenCLKernelArgMetadataAttr>(
+ attribute.getValue())) {
+ emitOpenCLKernelArgMetadata(clArgMetadata, llvmFunc, moduleTranslation);
+ } else if (auto strAttr =
+ mlir::dyn_cast<mlir::StringAttr>(attribute.getValue())) {
llvmFunc->addFnAttr(llvmAttrName, strAttr.getValue());
+ }
return mlir::success();
}
+ void emitOpenCLKernelArgMetadata(
+ cir::OpenCLKernelArgMetadataAttr clArgMetadata, llvm::Function *llvmFunc,
+ mlir::LLVM::ModuleTranslation &moduleTranslation) const {
+ llvm::LLVMContext &llvmContext = moduleTranslation.getLLVMContext();
+
+ SmallVector<llvm::Metadata *, 8> addressQuals;
+
+ auto getOpenCLArgInfoAddressSpace = [](cir::LangAddressSpace addressSpace) {
+ switch (addressSpace) {
+ case cir::LangAddressSpace::Default:
+ case cir::LangAddressSpace::OffloadPrivate:
+ return 0u;
+ case cir::LangAddressSpace::OffloadGlobal:
+ return 1u;
+ case cir::LangAddressSpace::OffloadConstant:
+ return 2u;
+ case cir::LangAddressSpace::OffloadLocal:
+ return 3u;
+ case cir::LangAddressSpace::OffloadGeneric:
+ return 4u;
+ case cir::LangAddressSpace::OffloadGlobalDevice:
+ return 5u;
+ case cir::LangAddressSpace::OffloadGlobalHost:
+ return 6u;
+ }
+ llvm_unreachable("unknown CIR language address space");
+ };
+
+ for (cir::LangAddressSpaceAttr addressSpace :
+ clArgMetadata.getAddrSpace().getAsRange<cir::LangAddressSpaceAttr>()) {
+ addressQuals.push_back(
+ llvm::ConstantAsMetadata::get(llvm::ConstantInt::get(
+ llvm::IntegerType::get(llvmContext, 32),
+ getOpenCLArgInfoAddressSpace(addressSpace.getValue()))));
+ }
+
+ llvmFunc->setMetadata("kernel_arg_addr_space",
+ llvm::MDNode::get(llvmContext, addressQuals));
+ llvmFunc->setMetadata(
+ "kernel_arg_access_qual",
+ getStringArrayMetadataNode(llvmContext, clArgMetadata.getAccessQual()));
+ llvmFunc->setMetadata(
+ "kernel_arg_type",
+ getStringArrayMetadataNode(llvmContext, clArgMetadata.getType()));
+ llvmFunc->setMetadata(
+ "kernel_arg_base_type",
+ getStringArrayMetadataNode(llvmContext, clArgMetadata.getBaseType()));
+ llvmFunc->setMetadata(
+ "kernel_arg_type_qual",
+ getStringArrayMetadataNode(llvmContext, clArgMetadata.getTypeQual()));
+ if (clArgMetadata.getName())
+ llvmFunc->setMetadata(
+ "kernel_arg_name",
+ getStringArrayMetadataNode(llvmContext, clArgMetadata.getName()));
+ }
+
+ llvm::MDNode *getStringArrayMetadataNode(llvm::LLVMContext &llvmContext,
+ mlir::ArrayAttr attrs) const {
+ SmallVector<llvm::Metadata *, 8> metadata;
+ for (mlir::StringAttr attr : attrs.getAsRange<mlir::StringAttr>())
+ metadata.push_back(llvm::MDString::get(llvmContext, attr.getValue()));
+ return llvm::MDNode::get(llvmContext, metadata);
+ }
+
// Translate CIR's module attributes to LLVM's module metadata
mlir::LogicalResult
amendModule(mlir::ModuleOp mod, mlir::NamedAttribute attribute,
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 e18a125098f64..f4823b61966cf 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 @@
// 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,9 +15,17 @@ 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-LABEL: define{{.*}} void @spir_addr_space_kernel_args
+// LLVM-SAME: !kernel_arg_addr_space ![[ADDR_SPACES:[0-9]+]]
+
kernel void global_device_host_kernel_args(
__attribute__((opencl_global_device)) int *D,
__attribute__((opencl_global_host)) int *H) {}
// 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-LABEL: define{{.*}} void @global_device_host_kernel_args
+// LLVM-SAME: !kernel_arg_addr_space ![[GLOBAL_DEVICE_HOST_ADDR_SPACES:[0-9]+]]
+// LLVM-DAG: ![[ADDR_SPACES]] = !{i32 1, i32 2, i32 3}
+// LLVM-DAG: ![[GLOBAL_DEVICE_HOST_ADDR_SPACES]] = !{i32 5, i32 6}
diff --git a/clang/test/CIR/CodeGenOpenCL/kernel-arg-info.cl b/clang/test/CIR/CodeGenOpenCL/kernel-arg-info.cl
index 7788195157715..8098843228226 100644
--- a/clang/test/CIR/CodeGenOpenCL/kernel-arg-info.cl
+++ b/clang/test/CIR/CodeGenOpenCL/kernel-arg-info.cl
@@ -4,6 +4,15 @@
// 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: %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-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: %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-ARGINFO
+
kernel void global_qualifier_kernel_args(
global int *globalintp, global int *restrict globalintrestrictp,
global const int *globalconstintp,
@@ -29,6 +38,14 @@ 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-ARGINFO-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]+]] !kernel_arg_name ![[GLOBAL_ARG_NAMES:[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: ![[GLOBAL_ARG_NAMES]] = !{!"globalintp", !"globalintrestrictp", !"globalconstintp", !"globalconstintrestrictp", !"globalconstvolatileintp", !"globalconstvolatileintrestrictp", !"globalvolatileintp", !"globalvolatileintrestrictp"}
+
kernel void constant_kernel_args(constant int *constantintp,
constant int *restrict constantintrestrictp) {}
@@ -48,6 +65,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-ARGINFO-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]+]] !kernel_arg_name ![[CONSTANT_ARG_NAMES:[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: ![[CONSTANT_ARG_NAMES]] = !{!"constantintp", !"constantintrestrictp"}
+
kernel void local_qualifier_kernel_args(
local int *localintp, local int *restrict localintrestrictp,
local const int *localconstintp,
@@ -73,6 +98,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-ARGINFO-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]] !kernel_arg_name ![[LOCAL_ARG_NAMES:[0-9]+]]
+// LLVM-DAG: ![[LOCAL_ADDR_SPACES]] = !{i32 3, i32 3, i32 3, i32 3, i32 3, i32 3, i32 3, i32 3}
+// 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 +123,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-ARGINFO-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]+]] !kernel_arg_name ![[PRIVATE_ARG_NAMES:[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: ![[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 +151,15 @@ 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-ARGINFO-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]+]] !kernel_arg_name ![[TYPEDEF_ARG_NAMES:[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: ![[TYPEDEF_ARG_NAMES]] = !{!"X", !"Y"}
+
typedef char char16 __attribute__((ext_vector_type(16)));
__kernel void vector_typedef_kernel_arg(__global char16 arg[]) {}
@@ -132,6 +179,11 @@ __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_type ![[VECTOR_TYPEDEF_ARG_TYPES:[0-9]+]]
+// LLVM-ARGINFO-DAG: define{{.*}} void @vector_typedef_kernel_arg{{.+}} !kernel_arg_name ![[VECTOR_TYPEDEF_ARG_NAMES:[0-9]+]]
+// LLVM-DAG: ![[VECTOR_TYPEDEF_ARG_TYPES]] = !{!"char16*"}
+// LLVM-ARGINFO-DAG: ![[VECTOR_TYPEDEF_ARG_NAMES]] = !{!"arg"}
+
kernel void signed_char_kernel_args(signed char sc1,
global const signed char *sc2) {}
@@ -150,3 +202,11 @@ 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-ARGINFO-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]+]] !kernel_arg_name ![[SIGNED_CHAR_ARG_NAMES:[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: ![[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 dd90ac27d6ec5..85c300d3b005b 100644
--- a/clang/test/CIR/CodeGenOpenCL/kernel-arg-metadata.cl
+++ b/clang/test/CIR/CodeGenOpenCL/kernel-arg-metadata.cl
@@ -1,7 +1,19 @@
// 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
__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]] = !{}
More information about the llvm-branch-commits
mailing list