https://github.com/steffenlarsen updated https://github.com/llvm/llvm-project/pull/226996
>From cd71c07d6f88cf4828fd7eed7946e0906e63a1ef Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Mon, 28 Sep 2026 09:24:12 -0500 Subject: [PATCH 1/4] [CIR] Fix address space of string literals This commit fixes the address space of string literals in CIR. The address space of string literals is determined by the AST type, which is always the default address space for address space agnostic languages. The global variable that holds the string literal may live in a different address space, e.g. OpenCL constant address space or HIP device address space. Therefore, we need to cast the global variable to the pointer type of the AST type when emitting a reference to the string literal. Assisted-by: Claude Opus 5.5 Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/lib/CIR/CodeGen/CIRGenExpr.cpp | 5 + clang/lib/CIR/CodeGen/CIRGenModule.cpp | 27 +++++- clang/lib/CIR/CodeGen/CIRGenModule.h | 7 ++ .../CodeGenHIP/string-literal-addrspace.hip | 97 +++++++++++++++++++ 4 files changed, 131 insertions(+), 5 deletions(-) create mode 100644 clang/test/CIR/CodeGenHIP/string-literal-addrspace.hip diff --git a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp index 7688fcc3cc337..6d239ae16327f 100644 --- a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp @@ -1585,6 +1585,11 @@ LValue CIRGenFunction::emitStringLiteralLValue(const StringLiteral *e, unsigned align = *(globalOp.getAlignment()); mlir::Value addr = builder.createGetGlobal(getLoc(e->getSourceRange()), globalOp); + cir::PointerType destPtrTy = + builder.getPointerTo(globalOp.getSymType(), + cgm.getTypes().getPointerAddressSpace(e->getType())); + if (addr.getType() != destPtrTy) + addr = performAddrSpaceCast(addr, destPtrTy); return makeAddrLValue( Address(addr, globalOp.getSymType(), CharUnits::fromQuantity(align)), e->getType(), AlignmentSource::Decl); diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index 007788d7e27c7..d09a3504e284a 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -2258,12 +2258,14 @@ static cir::GlobalOp generateStringLiteral(mlir::Location loc, mlir::TypedAttr c, cir::GlobalLinkageKind lt, CIRGenModule &cgm, StringRef globalName, CharUnits alignment) { - assert(!cir::MissingFeatures::addressSpace()); + mlir::ptr::MemorySpaceAttrInterface addrSpace = cir::toCIRAddressSpaceAttr( + cgm.getMLIRContext(), cgm.getGlobalConstantAddressSpace()); // Create a global variable for this string // FIXME(cir): check for insertion point in module level. - cir::GlobalOp gv = cgm.createGlobalOp(loc, globalName, c.getType(), - !cgm.getLangOpts().WritableStrings); + cir::GlobalOp gv = + cgm.createGlobalOp(loc, globalName, c.getType(), + !cgm.getLangOpts().WritableStrings, addrSpace); // Set up extra information and add to the module gv.setAlignmentAttr(cgm.getSize(alignment)); @@ -2356,12 +2358,27 @@ CIRGenModule::getAddrOfConstantStringFromLiteral(const StringLiteral *s, cir::GlobalOp gv = getGlobalForStringLiteral(s, name); auto arrayTy = mlir::dyn_cast<cir::ArrayType>(gv.getSymType()); assert(arrayTy && "String literal must be array"); - assert(!cir::MissingFeatures::addressSpace()); - cir::PointerType ptrTy = getBuilder().getPointerTo(arrayTy.getElementType()); + cir::PointerType ptrTy = getBuilder().getPointerTo( + arrayTy.getElementType(), + getTypes().getPointerAddressSpace(s->getType())); return builder.getGlobalViewAttr(ptrTy, gv); } +LangAS CIRGenModule::getGlobalConstantAddressSpace() const { + if (langOpts.OpenCL) + return LangAS::opencl_constant; + if (langOpts.SYCLIsDevice) { + errorNYI("SYCL global constant address space"); + return LangAS::Default; + } + if (langOpts.HIP && langOpts.CUDAIsDevice && getTriple().isSPIRV()) + return LangAS::cuda_device; + if (std::optional<LangAS> constAS = getTarget().getConstantAddressSpace()) + return *constAS; + return LangAS::Default; +} + // TODO(cir): this could be a common AST helper for both CIR and LLVM codegen. LangAS CIRGenModule::getLangTempAllocaAddressSpace() const { if (getLangOpts().OpenCL) diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h index 3fb95f346536d..208ba69dfa65f 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.h +++ b/clang/lib/CIR/CodeGen/CIRGenModule.h @@ -460,6 +460,13 @@ class CIRGenModule : public CIRGenTypeCache { getAddrOfConstantStringFromLiteral(const StringLiteral *s, llvm::StringRef name = ".str"); + /// Return the AST address space of constant literal, which is used to emit + /// the constant literal as global variable in CIR. This is not necessarily + /// the address space of the constant literal in AST. For address space + /// agnostic language, e.g. C++, constant literal in AST is always in the + /// default address space. + LangAS getGlobalConstantAddressSpace() const; + /// Returns the address space for temporary allocations in the language. This /// ensures that the allocated variable's address space matches the /// expectations of the AST, rather than using the target's allocation address diff --git a/clang/test/CIR/CodeGenHIP/string-literal-addrspace.hip b/clang/test/CIR/CodeGenHIP/string-literal-addrspace.hip new file mode 100644 index 0000000000000..a9d060a97b3eb --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/string-literal-addrspace.hip @@ -0,0 +1,97 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fcuda-is-device -fclangir -emit-cir %s -o - \ +// RUN: | FileCheck --check-prefix=CIR-SPV %s +// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fcuda-is-device -fclangir -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=LLVM-SPV %s +// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fcuda-is-device -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=OGCG-SPV %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device -fclangir -emit-cir %s -o - \ +// RUN: | FileCheck --check-prefix=CIR-GCN %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device -fclangir -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=LLVM-GCN %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=OGCG-GCN %s + +// String literals are emitted in the target's constant address space for +// globals (CrossWorkgroup on SPIR-V, constant on AMDGCN) and are referenced +// through an addrspacecast to the default address space. + +#define __device__ __attribute__((device)) + +__device__ void take(const char *); + +// CIR-SPV: cir.global "private" constant cir_private dso_local target_address_space(1) @".str" = #cir.const_array<"hi" : !cir.array<!s8i x 2>, trailing_zeros> : !cir.array<!s8i x 3> +// CIR-SPV: cir.global "private" constant cir_private dso_local target_address_space(1) @".str.1" = #cir.const_array<"glob" : !cir.array<!s8i x 4>, trailing_zeros> : !cir.array<!s8i x 5> +// CIR-SPV: cir.global external target_address_space(1) @gp = #cir.global_view<@".str.1"> : !cir.ptr<!s8i, target_address_space(4)> + +// CIR-GCN: cir.global "private" constant cir_private dso_local target_address_space(4) @".str" = #cir.const_array<"hi" : !cir.array<!s8i x 2>, trailing_zeros> : !cir.array<!s8i x 3> +// CIR-GCN: cir.global "private" constant cir_private dso_local target_address_space(4) @".str.1" = #cir.const_array<"glob" : !cir.array<!s8i x 4>, trailing_zeros> : !cir.array<!s8i x 5> +// CIR-GCN: cir.global external target_address_space(1) @gp = #cir.global_view<@".str.1"> : !cir.ptr<!s8i> + +// LLVM-SPV: @.str = private addrspace(1) constant [3 x i8] c"hi\00" +// LLVM-SPV: @.str.1 = private addrspace(1) constant [5 x i8] c"glob\00" +// LLVM-SPV: @gp = addrspace(1) externally_initialized global ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str.1 to ptr addrspace(4)) + +// OGCG-SPV: @.str = private unnamed_addr addrspace(1) constant [3 x i8] c"hi\00" +// OGCG-SPV: @.str.1 = private unnamed_addr addrspace(1) constant [5 x i8] c"glob\00" +// OGCG-SPV: @gp = addrspace(1) externally_initialized global ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str.1 to ptr addrspace(4)) + +// LLVM-GCN: @.str = private addrspace(4) constant [3 x i8] c"hi\00" +// LLVM-GCN: @.str.1 = private addrspace(4) constant [5 x i8] c"glob\00" +// LLVM-GCN: @gp = addrspace(1) externally_initialized global ptr addrspacecast (ptr addrspace(4) @.str.1 to ptr) + +// OGCG-GCN: @.str = private unnamed_addr addrspace(4) constant [3 x i8] c"hi\00" +// OGCG-GCN: @.str.1 = private unnamed_addr addrspace(4) constant [5 x i8] c"glob\00" +// OGCG-GCN: @gp = addrspace(1) externally_initialized global ptr addrspacecast (ptr addrspace(4) @.str.1 to ptr) + +// CIR-SPV-LABEL: cir.func{{.*}} @_Z9call_takev +// CIR-SPV: %[[G:.*]] = cir.get_global @".str" : !cir.ptr<!cir.array<!s8i x 3>, target_address_space(1)> +// CIR-SPV: %[[C:.*]] = cir.cast address_space %[[G]] : !cir.ptr<!cir.array<!s8i x 3>, target_address_space(1)> -> !cir.ptr<!cir.array<!s8i x 3>, target_address_space(4)> +// CIR-SPV: %[[D:.*]] = cir.cast array_to_ptrdecay %[[C]] : !cir.ptr<!cir.array<!s8i x 3>, target_address_space(4)> -> !cir.ptr<!s8i, target_address_space(4)> +// CIR-SPV: cir.call @_Z4takePKc(%[[D]]) + +// CIR-GCN-LABEL: cir.func{{.*}} @_Z9call_takev +// CIR-GCN: %[[G:.*]] = cir.get_global @".str" : !cir.ptr<!cir.array<!s8i x 3>, target_address_space(4)> +// CIR-GCN: %[[C:.*]] = cir.cast address_space %[[G]] : !cir.ptr<!cir.array<!s8i x 3>, target_address_space(4)> -> !cir.ptr<!cir.array<!s8i x 3>> +// CIR-GCN: %[[D:.*]] = cir.cast array_to_ptrdecay %[[C]] : !cir.ptr<!cir.array<!s8i x 3>> -> !cir.ptr<!s8i> +// CIR-GCN: cir.call @_Z4takePKc(%[[D]]) + +// LLVM-SPV-LABEL: define{{.*}} void @_Z9call_takev +// LLVM-SPV: call{{.*}} void @_Z4takePKc(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @.str to ptr addrspace(4))) + +// OGCG-SPV-LABEL: define{{.*}} void @_Z9call_takev +// OGCG-SPV: call{{.*}} void @_Z4takePKc(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @.str to ptr addrspace(4))) + +// LLVM-GCN-LABEL: define{{.*}} void @_Z9call_takev +// LLVM-GCN: call void @_Z4takePKc(ptr noundef addrspacecast (ptr addrspace(4) @.str to ptr)) + +// OGCG-GCN-LABEL: define{{.*}} void @_Z9call_takev +// OGCG-GCN: call void @_Z4takePKc(ptr noundef addrspacecast (ptr addrspace(4) @.str to ptr)) +__device__ void call_take() { take("hi"); } + +// A second use of the same literal reuses the cached global. + +// CIR-SPV-LABEL: cir.func{{.*}} @_Z3retv +// CIR-SPV: %[[G:.*]] = cir.get_global @".str" : !cir.ptr<!cir.array<!s8i x 3>, target_address_space(1)> +// CIR-SPV: cir.cast address_space %[[G]] : !cir.ptr<!cir.array<!s8i x 3>, target_address_space(1)> -> !cir.ptr<!cir.array<!s8i x 3>, target_address_space(4)> + +// CIR-GCN-LABEL: cir.func{{.*}} @_Z3retv +// CIR-GCN: %[[G:.*]] = cir.get_global @".str" : !cir.ptr<!cir.array<!s8i x 3>, target_address_space(4)> +// CIR-GCN: cir.cast address_space %[[G]] : !cir.ptr<!cir.array<!s8i x 3>, target_address_space(4)> -> !cir.ptr<!cir.array<!s8i x 3>> + +// LLVM-SPV-LABEL: define{{.*}} ptr addrspace(4) @_Z3retv +// LLVM-SPV: store ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str to ptr addrspace(4)) + +// OGCG-SPV-LABEL: define{{.*}} ptr addrspace(4) @_Z3retv +// OGCG-SPV: ret ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str to ptr addrspace(4)) + +// LLVM-GCN-LABEL: define{{.*}} ptr @_Z3retv +// LLVM-GCN: store ptr addrspacecast (ptr addrspace(4) @.str to ptr) + +// OGCG-GCN-LABEL: define{{.*}} ptr @_Z3retv +// OGCG-GCN: ret ptr addrspacecast (ptr addrspace(4) @.str to ptr) +__device__ const char *ret() { return "hi"; } + +// Initializing a global goes through constant emission instead. +__device__ const char *gp = "glob"; >From ddeeecbacee5b57697a1d0df5f2d9a3ac3d2d255 Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Mon, 28 Sep 2026 11:32:47 -0500 Subject: [PATCH 2/4] Move implementation to common util but keep wrapper for now Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/include/clang/CodeGenUtils/CodeGenUtils.h | 8 ++++++++ clang/lib/CIR/CodeGen/CIRGenModule.cpp | 13 +++++-------- clang/lib/CodeGen/CodeGenModule.cpp | 17 +---------------- clang/lib/CodeGenUtils/CodeGenUtils.cpp | 14 ++++++++++++++ 4 files changed, 28 insertions(+), 24 deletions(-) diff --git a/clang/include/clang/CodeGenUtils/CodeGenUtils.h b/clang/include/clang/CodeGenUtils/CodeGenUtils.h index b24f457f9537c..b2cc416330eb3 100644 --- a/clang/include/clang/CodeGenUtils/CodeGenUtils.h +++ b/clang/include/clang/CodeGenUtils/CodeGenUtils.h @@ -44,6 +44,14 @@ bool hasUnwindExceptions(const LangOptions &LangOpts); /// Helper method to check if the underlying ABI is AAPCS bool isAAPCS(const TargetInfo &TargetInfo); +/// Return the AST address space of constant literal, which is used to emit +/// the constant literal as global variable in LLVM IR. +/// Note: This is not necessarily the address space of the constant literal +/// in AST. For address space agnostic language, e.g. C++, constant literal +/// in AST is always in default address space. +LangAS getGlobalConstantAddressSpace(const LangOptions &LangOpts, + const TargetInfo &Target); + bool isInitializerOfDynamicClass(const CXXCtorInitializer *BaseInit); /// Check that a call to a target-specific builtin has the required target diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index d09a3504e284a..8700e7cbd0fd1 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -2366,17 +2366,14 @@ CIRGenModule::getAddrOfConstantStringFromLiteral(const StringLiteral *s, } LangAS CIRGenModule::getGlobalConstantAddressSpace() const { - if (langOpts.OpenCL) - return LangAS::opencl_constant; - if (langOpts.SYCLIsDevice) { + LangAS as = + CodeGenUtils::getGlobalConstantAddressSpace(langOpts, getTarget()); + // CIR cannot represent SYCL address spaces yet. + if (as == LangAS::sycl_global) { errorNYI("SYCL global constant address space"); return LangAS::Default; } - if (langOpts.HIP && langOpts.CUDAIsDevice && getTriple().isSPIRV()) - return LangAS::cuda_device; - if (std::optional<LangAS> constAS = getTarget().getConstantAddressSpace()) - return *constAS; - return LangAS::Default; + return as; } // TODO(cir): this could be a common AST helper for both CIR and LLVM codegen. diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 7274a8588670f..bb24841c522b1 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -6446,22 +6446,7 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) { } LangAS CodeGenModule::GetGlobalConstantAddressSpace() const { - // OpenCL v1.2 s6.5.3: a string literal is in the constant address space. - if (LangOpts.OpenCL) - return LangAS::opencl_constant; - if (LangOpts.SYCLIsDevice) - return LangAS::sycl_global; - if (LangOpts.HIP && LangOpts.CUDAIsDevice && getTriple().isSPIRV()) - // For HIPSPV map literals to cuda_device (maps to CrossWorkGroup in SPIR-V) - // instead of default AS (maps to Generic in SPIR-V). Otherwise, we end up - // with OpVariable instructions with Generic storage class which is not - // allowed (SPIR-V V1.6 s3.42.8). Also, mapping literals to SPIR-V - // UniformConstant storage class is not viable as pointers to it may not be - // casted to Generic pointers which are used to model HIP's "flat" pointers. - return LangAS::cuda_device; - if (auto AS = getTarget().getConstantAddressSpace()) - return *AS; - return LangAS::Default; + return CodeGenUtils::getGlobalConstantAddressSpace(LangOpts, getTarget()); } // In address space agnostic languages, string literals are in default address diff --git a/clang/lib/CodeGenUtils/CodeGenUtils.cpp b/clang/lib/CodeGenUtils/CodeGenUtils.cpp index 4fd78d6997e95..9f625a0089539 100644 --- a/clang/lib/CodeGenUtils/CodeGenUtils.cpp +++ b/clang/lib/CodeGenUtils/CodeGenUtils.cpp @@ -113,6 +113,20 @@ bool hasUnwindExceptions(const LangOptions &LangOpts) { bool isAAPCS(const TargetInfo &TargetInfo) { return TargetInfo.getABI().starts_with("aapcs"); } + +LangAS getGlobalConstantAddressSpace(const LangOptions &LangOpts, + const TargetInfo &Target) { + if (LangOpts.OpenCL) + return LangAS::opencl_constant; + if (LangOpts.SYCLIsDevice) + return LangAS::sycl_global; + if (LangOpts.HIP && LangOpts.CUDAIsDevice && Target.getTriple().isSPIRV()) + return LangAS::cuda_device; + if (auto AS = Target.getConstantAddressSpace()) + return *AS; + return LangAS::Default; +} + bool isInitializerOfDynamicClass(const CXXCtorInitializer *BaseInit) { const Type *BaseType = BaseInit->getBaseClass(); return BaseType->castAsCXXRecordDecl()->isDynamicClass(); >From 69f313b79cfa7ced1bde0585cec756282d587dec Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Mon, 28 Sep 2026 11:36:02 -0500 Subject: [PATCH 3/4] Compare addr space instead of type Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/lib/CIR/CodeGen/CIRGenExpr.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp index 6d239ae16327f..56c18bc1bb8d3 100644 --- a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp @@ -1585,11 +1585,11 @@ LValue CIRGenFunction::emitStringLiteralLValue(const StringLiteral *e, unsigned align = *(globalOp.getAlignment()); mlir::Value addr = builder.createGetGlobal(getLoc(e->getSourceRange()), globalOp); - cir::PointerType destPtrTy = - builder.getPointerTo(globalOp.getSymType(), - cgm.getTypes().getPointerAddressSpace(e->getType())); - if (addr.getType() != destPtrTy) - addr = performAddrSpaceCast(addr, destPtrTy); + mlir::ptr::MemorySpaceAttrInterface destAS = + cgm.getTypes().getPointerAddressSpace(e->getType()); + if (mlir::cast<cir::PointerType>(addr.getType()).getAddrSpace() != destAS) + addr = performAddrSpaceCast( + addr, builder.getPointerTo(globalOp.getSymType(), destAS)); return makeAddrLValue( Address(addr, globalOp.getSymType(), CharUnits::fromQuantity(align)), e->getType(), AlignmentSource::Decl); >From 0ac6bd7e29821d548b5e70c1255a3933a719053c Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Tue, 29 Sep 2026 00:26:12 -0500 Subject: [PATCH 4/4] Add more testing Signed-off-by: Steffen Holst Larsen <[email protected]> --- .../CodeGenHIP/string-literal-addrspace.hip | 129 ++++++++++++++++++ .../CodeGenOpenCL/string-literal-addrspace.cl | 111 +++++++++++++++ .../string-literal-addrspace-nyi.cpp | 20 +++ 3 files changed, 260 insertions(+) create mode 100644 clang/test/CIR/CodeGenOpenCL/string-literal-addrspace.cl create mode 100644 clang/test/CIR/CodeGenSYCL/string-literal-addrspace-nyi.cpp diff --git a/clang/test/CIR/CodeGenHIP/string-literal-addrspace.hip b/clang/test/CIR/CodeGenHIP/string-literal-addrspace.hip index a9d060a97b3eb..285c2e7325431 100644 --- a/clang/test/CIR/CodeGenHIP/string-literal-addrspace.hip +++ b/clang/test/CIR/CodeGenHIP/string-literal-addrspace.hip @@ -21,29 +21,62 @@ __device__ void take(const char *); +// CIR emits the static local before the string literals; classic codegen +// emits it in source order. + +// CIR-SPV: cir.global "private" internal dso_local target_address_space(1) @_ZZ12static_localvE1p = #cir.global_view<@".str.3"> : !cir.ptr<!s8i, target_address_space(4)> // CIR-SPV: cir.global "private" constant cir_private dso_local target_address_space(1) @".str" = #cir.const_array<"hi" : !cir.array<!s8i x 2>, trailing_zeros> : !cir.array<!s8i x 3> // CIR-SPV: cir.global "private" constant cir_private dso_local target_address_space(1) @".str.1" = #cir.const_array<"glob" : !cir.array<!s8i x 4>, trailing_zeros> : !cir.array<!s8i x 5> // CIR-SPV: cir.global external target_address_space(1) @gp = #cir.global_view<@".str.1"> : !cir.ptr<!s8i, target_address_space(4)> +// CIR-SPV: cir.global "private" constant cir_private dso_local target_address_space(1) @__func__._Z9func_namev = #cir.const_array<"func_name" : !cir.array<!s8i x 9>, trailing_zeros> : !cir.array<!s8i x 10> +// CIR-SPV: cir.global "private" constant cir_private dso_local target_address_space(1) @".str.2" = #cir.const_array<"abc" : !cir.array<!s8i x 3>, trailing_zeros> : !cir.array<!s8i x 4> +// CIR-SPV: cir.global "private" constant cir_private dso_local target_address_space(1) @".str.3" = #cir.const_array<"st" : !cir.array<!s8i x 2>, trailing_zeros> : !cir.array<!s8i x 3> +// CIR-SPV: cir.global "private" constant cir_private dso_local target_address_space(1) @".str.4" = #cir.const_array<[#cir.int<119> : !s32i], trailing_zeros> : !cir.array<!s32i x 2> +// CIR-GCN: cir.global "private" internal dso_local target_address_space(1) @_ZZ12static_localvE1p = #cir.global_view<@".str.3"> : !cir.ptr<!s8i> // CIR-GCN: cir.global "private" constant cir_private dso_local target_address_space(4) @".str" = #cir.const_array<"hi" : !cir.array<!s8i x 2>, trailing_zeros> : !cir.array<!s8i x 3> // CIR-GCN: cir.global "private" constant cir_private dso_local target_address_space(4) @".str.1" = #cir.const_array<"glob" : !cir.array<!s8i x 4>, trailing_zeros> : !cir.array<!s8i x 5> // CIR-GCN: cir.global external target_address_space(1) @gp = #cir.global_view<@".str.1"> : !cir.ptr<!s8i> +// CIR-GCN: cir.global "private" constant cir_private dso_local target_address_space(4) @__func__._Z9func_namev = #cir.const_array<"func_name" : !cir.array<!s8i x 9>, trailing_zeros> : !cir.array<!s8i x 10> +// CIR-GCN: cir.global "private" constant cir_private dso_local target_address_space(4) @".str.2" = #cir.const_array<"abc" : !cir.array<!s8i x 3>, trailing_zeros> : !cir.array<!s8i x 4> +// CIR-GCN: cir.global "private" constant cir_private dso_local target_address_space(4) @".str.3" = #cir.const_array<"st" : !cir.array<!s8i x 2>, trailing_zeros> : !cir.array<!s8i x 3> +// CIR-GCN: cir.global "private" constant cir_private dso_local target_address_space(4) @".str.4" = #cir.const_array<[#cir.int<119> : !s32i], trailing_zeros> : !cir.array<!s32i x 2> +// LLVM-SPV: @_ZZ12static_localvE1p = internal addrspace(1) global ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str.3 to ptr addrspace(4)) // LLVM-SPV: @.str = private addrspace(1) constant [3 x i8] c"hi\00" // LLVM-SPV: @.str.1 = private addrspace(1) constant [5 x i8] c"glob\00" // LLVM-SPV: @gp = addrspace(1) externally_initialized global ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str.1 to ptr addrspace(4)) +// LLVM-SPV: @__func__._Z9func_namev = private addrspace(1) constant [10 x i8] c"func_name\00" +// LLVM-SPV: @.str.2 = private addrspace(1) constant [4 x i8] c"abc\00" +// LLVM-SPV: @.str.3 = private addrspace(1) constant [3 x i8] c"st\00" +// LLVM-SPV: @.str.4 = private addrspace(1) constant [2 x i32] [i32 119, i32 0] // OGCG-SPV: @.str = private unnamed_addr addrspace(1) constant [3 x i8] c"hi\00" // OGCG-SPV: @.str.1 = private unnamed_addr addrspace(1) constant [5 x i8] c"glob\00" // OGCG-SPV: @gp = addrspace(1) externally_initialized global ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str.1 to ptr addrspace(4)) +// OGCG-SPV: @__func__._Z9func_namev = private unnamed_addr addrspace(1) constant [10 x i8] c"func_name\00" +// OGCG-SPV: @.str.2 = private unnamed_addr addrspace(1) constant [4 x i8] c"abc\00" +// OGCG-SPV: @_ZZ12static_localvE1p = internal addrspace(1) global ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str.3 to ptr addrspace(4)) +// OGCG-SPV: @.str.3 = private unnamed_addr addrspace(1) constant [3 x i8] c"st\00" +// OGCG-SPV: @.str.4 = private unnamed_addr addrspace(1) constant [2 x i32] [i32 119, i32 0] +// LLVM-GCN: @_ZZ12static_localvE1p = internal addrspace(1) global ptr addrspacecast (ptr addrspace(4) @.str.3 to ptr) // LLVM-GCN: @.str = private addrspace(4) constant [3 x i8] c"hi\00" // LLVM-GCN: @.str.1 = private addrspace(4) constant [5 x i8] c"glob\00" // LLVM-GCN: @gp = addrspace(1) externally_initialized global ptr addrspacecast (ptr addrspace(4) @.str.1 to ptr) +// LLVM-GCN: @__func__._Z9func_namev = private addrspace(4) constant [10 x i8] c"func_name\00" +// LLVM-GCN: @.str.2 = private addrspace(4) constant [4 x i8] c"abc\00" +// LLVM-GCN: @.str.3 = private addrspace(4) constant [3 x i8] c"st\00" +// LLVM-GCN: @.str.4 = private addrspace(4) constant [2 x i32] [i32 119, i32 0] // OGCG-GCN: @.str = private unnamed_addr addrspace(4) constant [3 x i8] c"hi\00" // OGCG-GCN: @.str.1 = private unnamed_addr addrspace(4) constant [5 x i8] c"glob\00" // OGCG-GCN: @gp = addrspace(1) externally_initialized global ptr addrspacecast (ptr addrspace(4) @.str.1 to ptr) +// OGCG-GCN: @__func__._Z9func_namev = private unnamed_addr addrspace(4) constant [10 x i8] c"func_name\00" +// OGCG-GCN: @.str.2 = private unnamed_addr addrspace(4) constant [4 x i8] c"abc\00" +// OGCG-GCN: @_ZZ12static_localvE1p = internal addrspace(1) global ptr addrspacecast (ptr addrspace(4) @.str.3 to ptr) +// OGCG-GCN: @.str.3 = private unnamed_addr addrspace(4) constant [3 x i8] c"st\00" +// OGCG-GCN: @.str.4 = private unnamed_addr addrspace(4) constant [2 x i32] [i32 119, i32 0] // CIR-SPV-LABEL: cir.func{{.*}} @_Z9call_takev // CIR-SPV: %[[G:.*]] = cir.get_global @".str" : !cir.ptr<!cir.array<!s8i x 3>, target_address_space(1)> @@ -95,3 +128,99 @@ __device__ const char *ret() { return "hi"; } // Initializing a global goes through constant emission instead. __device__ const char *gp = "glob"; + +// __func__ is emitted like a string literal. + +// CIR-SPV-LABEL: cir.func{{.*}} @_Z9func_namev +// CIR-SPV: %[[G:.*]] = cir.get_global @__func__._Z9func_namev : !cir.ptr<!cir.array<!s8i x 10>, target_address_space(1)> +// CIR-SPV: cir.cast address_space %[[G]] : !cir.ptr<!cir.array<!s8i x 10>, target_address_space(1)> -> !cir.ptr<!cir.array<!s8i x 10>, target_address_space(4)> + +// CIR-GCN-LABEL: cir.func{{.*}} @_Z9func_namev +// CIR-GCN: %[[G:.*]] = cir.get_global @__func__._Z9func_namev : !cir.ptr<!cir.array<!s8i x 10>, target_address_space(4)> +// CIR-GCN: cir.cast address_space %[[G]] : !cir.ptr<!cir.array<!s8i x 10>, target_address_space(4)> -> !cir.ptr<!cir.array<!s8i x 10>> + +// LLVM-SPV-LABEL: define{{.*}} void @_Z9func_namev +// LLVM-SPV: call{{.*}} void @_Z4takePKc(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @__func__._Z9func_namev to ptr addrspace(4))) + +// OGCG-SPV-LABEL: define{{.*}} void @_Z9func_namev +// OGCG-SPV: call{{.*}} void @_Z4takePKc(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @__func__._Z9func_namev to ptr addrspace(4))) + +// LLVM-GCN-LABEL: define{{.*}} void @_Z9func_namev +// LLVM-GCN: call void @_Z4takePKc(ptr noundef addrspacecast (ptr addrspace(4) @__func__._Z9func_namev to ptr)) + +// OGCG-GCN-LABEL: define{{.*}} void @_Z9func_namev +// OGCG-GCN: call void @_Z4takePKc(ptr noundef addrspacecast (ptr addrspace(4) @__func__._Z9func_namev to ptr)) +__device__ void func_name() { take(__func__); } + +// Subscripting a literal indexes through the cast pointer. + +// CIR-SPV-LABEL: cir.func{{.*}} @_Z9subscripti +// CIR-SPV: %[[G:.*]] = cir.get_global @".str.2" : !cir.ptr<!cir.array<!s8i x 4>, target_address_space(1)> +// CIR-SPV: %[[C:.*]] = cir.cast address_space %[[G]] : !cir.ptr<!cir.array<!s8i x 4>, target_address_space(1)> -> !cir.ptr<!cir.array<!s8i x 4>, target_address_space(4)> +// CIR-SPV: cir.get_element %[[C]][{{.*}}] : !cir.ptr<!cir.array<!s8i x 4>, target_address_space(4)> -> !cir.ptr<!s8i, target_address_space(4)> + +// CIR-GCN-LABEL: cir.func{{.*}} @_Z9subscripti +// CIR-GCN: %[[G:.*]] = cir.get_global @".str.2" : !cir.ptr<!cir.array<!s8i x 4>, target_address_space(4)> +// CIR-GCN: %[[C:.*]] = cir.cast address_space %[[G]] : !cir.ptr<!cir.array<!s8i x 4>, target_address_space(4)> -> !cir.ptr<!cir.array<!s8i x 4>> +// CIR-GCN: cir.get_element %[[C]][{{.*}}] : !cir.ptr<!cir.array<!s8i x 4>> -> !cir.ptr<!s8i> + +// LLVM-SPV-LABEL: define{{.*}} i8 @_Z9subscripti +// LLVM-SPV: getelementptr {{.*}}[4 x i8], ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str.2 to ptr addrspace(4)) + +// OGCG-SPV-LABEL: define{{.*}} i8 @_Z9subscripti +// OGCG-SPV: getelementptr {{.*}}[4 x i8], ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str.2 to ptr addrspace(4)) + +// LLVM-GCN-LABEL: define{{.*}} i8 @_Z9subscripti +// LLVM-GCN: getelementptr {{.*}}[4 x i8], ptr addrspacecast (ptr addrspace(4) @.str.2 to ptr) + +// OGCG-GCN-LABEL: define{{.*}} i8 @_Z9subscripti +// OGCG-GCN: getelementptr {{.*}}[4 x i8], ptr addrspacecast (ptr addrspace(4) @.str.2 to ptr) +__device__ char subscript(int i) { return "abc"[i]; } + +// Binding a reference to a literal goes through constant emission and reuses +// the cached global. + +// CIR-SPV-LABEL: cir.func{{.*}} @_Z3refv +// CIR-SPV: cir.const #cir.global_view<@".str"> : !cir.ptr<!cir.array<!s8i x 3>, target_address_space(4)> + +// CIR-GCN-LABEL: cir.func{{.*}} @_Z3refv +// CIR-GCN: cir.const #cir.global_view<@".str"> : !cir.ptr<!cir.array<!s8i x 3>> + +// LLVM-SPV-LABEL: define{{.*}} ptr addrspace(4) @_Z3refv +// LLVM-SPV: store ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str to ptr addrspace(4)) + +// OGCG-SPV-LABEL: define{{.*}} ptr addrspace(4) @_Z3refv +// OGCG-SPV: store ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str to ptr addrspace(4)) + +// LLVM-GCN-LABEL: define{{.*}} ptr @_Z3refv +// LLVM-GCN: store ptr addrspacecast (ptr addrspace(4) @.str to ptr) + +// OGCG-GCN-LABEL: define{{.*}} ptr @_Z3refv +// OGCG-GCN: store ptr addrspacecast (ptr addrspace(4) @.str to ptr) +__device__ const char *ref() { const char (&r)[3] = "hi"; return r; } + +// A static local initializer goes through constant emission, like gp. +__device__ const char *static_local() { static const char *p = "st"; return p; } + +// Wide literals use the same path with a wider element type. + +// CIR-SPV-LABEL: cir.func{{.*}} @_Z4widev +// CIR-SPV: %[[G:.*]] = cir.get_global @".str.4" : !cir.ptr<!cir.array<!s32i x 2>, target_address_space(1)> +// CIR-SPV: cir.cast address_space %[[G]] : !cir.ptr<!cir.array<!s32i x 2>, target_address_space(1)> -> !cir.ptr<!cir.array<!s32i x 2>, target_address_space(4)> + +// CIR-GCN-LABEL: cir.func{{.*}} @_Z4widev +// CIR-GCN: %[[G:.*]] = cir.get_global @".str.4" : !cir.ptr<!cir.array<!s32i x 2>, target_address_space(4)> +// CIR-GCN: cir.cast address_space %[[G]] : !cir.ptr<!cir.array<!s32i x 2>, target_address_space(4)> -> !cir.ptr<!cir.array<!s32i x 2>> + +// LLVM-SPV-LABEL: define{{.*}} ptr addrspace(4) @_Z4widev +// LLVM-SPV: store ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str.4 to ptr addrspace(4)) + +// OGCG-SPV-LABEL: define{{.*}} ptr addrspace(4) @_Z4widev +// OGCG-SPV: ret ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str.4 to ptr addrspace(4)) + +// LLVM-GCN-LABEL: define{{.*}} ptr @_Z4widev +// LLVM-GCN: store ptr addrspacecast (ptr addrspace(4) @.str.4 to ptr) + +// OGCG-GCN-LABEL: define{{.*}} ptr @_Z4widev +// OGCG-GCN: ret ptr addrspacecast (ptr addrspace(4) @.str.4 to ptr) +__device__ const wchar_t *wide() { return L"w"; } diff --git a/clang/test/CIR/CodeGenOpenCL/string-literal-addrspace.cl b/clang/test/CIR/CodeGenOpenCL/string-literal-addrspace.cl new file mode 100644 index 0000000000000..6f337aaa4de6c --- /dev/null +++ b/clang/test/CIR/CodeGenOpenCL/string-literal-addrspace.cl @@ -0,0 +1,111 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -x cl -triple spirv64-unknown-unknown -cl-std=CL2.0 -O0 -fclangir -emit-cir %s -o - \ +// RUN: | FileCheck --check-prefix=CIR-SPV %s \ +// RUN: --implicit-check-not='cir.cast address_space' +// RUN: %clang_cc1 -x cl -triple spirv64-unknown-unknown -cl-std=CL2.0 -O0 -fclangir -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=LLVM-SPV %s +// RUN: %clang_cc1 -x cl -triple spirv64-unknown-unknown -cl-std=CL2.0 -O0 -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=OGCG-SPV %s + +// RUN: %clang_cc1 -x cl -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -O0 -fclangir -emit-cir %s -o - \ +// RUN: | FileCheck --check-prefix=CIR-GCN %s \ +// RUN: --implicit-check-not='cir.cast address_space' +// RUN: %clang_cc1 -x cl -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -O0 -fclangir -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=LLVM-GCN %s +// RUN: %clang_cc1 -x cl -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -O0 -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=OGCG-GCN %s + +// OpenCL string literals are __constant in the AST already, so the global is +// emitted in the constant address space and used without an address space +// cast. + +void take(const __constant char *); + +// CIR-SPV: cir.global "private" constant cir_private dso_local target_address_space(2) @".str" = #cir.const_array<"hi" : !cir.array<!s8i x 2>, trailing_zeros> : !cir.array<!s8i x 3> +// CIR-SPV: cir.global "private" constant cir_private dso_local target_address_space(2) @".str.1" = #cir.const_array<"abc" : !cir.array<!s8i x 3>, trailing_zeros> : !cir.array<!s8i x 4> +// CIR-SPV: cir.global "private" constant cir_private dso_local target_address_space(2) @__func__.func_name = #cir.const_array<"func_name" : !cir.array<!s8i x 9>, trailing_zeros> : !cir.array<!s8i x 10> + +// CIR-GCN: cir.global "private" constant cir_private dso_local target_address_space(4) @".str" = #cir.const_array<"hi" : !cir.array<!s8i x 2>, trailing_zeros> : !cir.array<!s8i x 3> +// CIR-GCN: cir.global "private" constant cir_private dso_local target_address_space(4) @".str.1" = #cir.const_array<"abc" : !cir.array<!s8i x 3>, trailing_zeros> : !cir.array<!s8i x 4> +// CIR-GCN: cir.global "private" constant cir_private dso_local target_address_space(4) @__func__.func_name = #cir.const_array<"func_name" : !cir.array<!s8i x 9>, trailing_zeros> : !cir.array<!s8i x 10> + +// LLVM-SPV: @.str = private addrspace(2) constant [3 x i8] c"hi\00" +// LLVM-SPV: @.str.1 = private addrspace(2) constant [4 x i8] c"abc\00" +// LLVM-SPV: @__func__.func_name = private addrspace(2) constant [10 x i8] c"func_name\00" + +// OGCG-SPV: @.str = private unnamed_addr addrspace(2) constant [3 x i8] c"hi\00" +// OGCG-SPV: @.str.1 = private unnamed_addr addrspace(2) constant [4 x i8] c"abc\00" +// OGCG-SPV: @__func__.func_name = private unnamed_addr addrspace(2) constant [10 x i8] c"func_name\00" + +// LLVM-GCN: @.str = private addrspace(4) constant [3 x i8] c"hi\00" +// LLVM-GCN: @.str.1 = private addrspace(4) constant [4 x i8] c"abc\00" +// LLVM-GCN: @__func__.func_name = private addrspace(4) constant [10 x i8] c"func_name\00" + +// OGCG-GCN: @.str = private unnamed_addr addrspace(4) constant [3 x i8] c"hi\00" +// OGCG-GCN: @.str.1 = private unnamed_addr addrspace(4) constant [4 x i8] c"abc\00" +// OGCG-GCN: @__func__.func_name = private unnamed_addr addrspace(4) constant [10 x i8] c"func_name\00" + +// CIR-SPV-LABEL: cir.func{{.*}} @call_take +// CIR-SPV: %[[G:.*]] = cir.get_global @".str" : !cir.ptr<!cir.array<!s8i x 3>, target_address_space(2)> +// CIR-SPV: %[[D:.*]] = cir.cast array_to_ptrdecay %[[G]] : !cir.ptr<!cir.array<!s8i x 3>, target_address_space(2)> -> !cir.ptr<!s8i, target_address_space(2)> +// CIR-SPV: cir.call @take(%[[D]]) + +// CIR-GCN-LABEL: cir.func{{.*}} @call_take +// CIR-GCN: %[[G:.*]] = cir.get_global @".str" : !cir.ptr<!cir.array<!s8i x 3>, target_address_space(4)> +// CIR-GCN: %[[D:.*]] = cir.cast array_to_ptrdecay %[[G]] : !cir.ptr<!cir.array<!s8i x 3>, target_address_space(4)> -> !cir.ptr<!s8i, target_address_space(4)> +// CIR-GCN: cir.call @take(%[[D]]) + +// LLVM-SPV-LABEL: define{{.*}} void @call_take +// LLVM-SPV: call{{.*}} void @take(ptr addrspace(2) noundef @.str) + +// OGCG-SPV-LABEL: define{{.*}} void @call_take +// OGCG-SPV: call{{.*}} void @take(ptr addrspace(2) noundef @.str) + +// LLVM-GCN-LABEL: define{{.*}} void @call_take +// LLVM-GCN: call void @take(ptr addrspace(4) noundef @.str) + +// OGCG-GCN-LABEL: define{{.*}} void @call_take +// OGCG-GCN: call void @take(ptr addrspace(4) noundef @.str) +void call_take(void) { take("hi"); } + +// CIR-SPV-LABEL: cir.func{{.*}} @subscript +// CIR-SPV: %[[G:.*]] = cir.get_global @".str.1" : !cir.ptr<!cir.array<!s8i x 4>, target_address_space(2)> +// CIR-SPV: cir.get_element %[[G]][{{.*}}] : !cir.ptr<!cir.array<!s8i x 4>, target_address_space(2)> -> !cir.ptr<!s8i, target_address_space(2)> + +// CIR-GCN-LABEL: cir.func{{.*}} @subscript +// CIR-GCN: %[[G:.*]] = cir.get_global @".str.1" : !cir.ptr<!cir.array<!s8i x 4>, target_address_space(4)> +// CIR-GCN: cir.get_element %[[G]][{{.*}}] : !cir.ptr<!cir.array<!s8i x 4>, target_address_space(4)> -> !cir.ptr<!s8i, target_address_space(4)> + +// LLVM-SPV-LABEL: define{{.*}} i8 @subscript +// LLVM-SPV: getelementptr {{.*}}[4 x i8], ptr addrspace(2) @.str.1 + +// OGCG-SPV-LABEL: define{{.*}} i8 @subscript +// OGCG-SPV: getelementptr {{.*}}[4 x i8], ptr addrspace(2) @.str.1 + +// LLVM-GCN-LABEL: define{{.*}} i8 @subscript +// LLVM-GCN: getelementptr {{.*}}[4 x i8], ptr addrspace(4) @.str.1 + +// OGCG-GCN-LABEL: define{{.*}} i8 @subscript +// OGCG-GCN: getelementptr {{.*}}[4 x i8], ptr addrspace(4) @.str.1 +char subscript(int i) { return "abc"[i]; } + +// CIR-SPV-LABEL: cir.func{{.*}} @func_name +// CIR-SPV: %[[G:.*]] = cir.get_global @__func__.func_name : !cir.ptr<!cir.array<!s8i x 10>, target_address_space(2)> +// CIR-SPV: cir.cast array_to_ptrdecay %[[G]] + +// CIR-GCN-LABEL: cir.func{{.*}} @func_name +// CIR-GCN: %[[G:.*]] = cir.get_global @__func__.func_name : !cir.ptr<!cir.array<!s8i x 10>, target_address_space(4)> +// CIR-GCN: cir.cast array_to_ptrdecay %[[G]] + +// LLVM-SPV-LABEL: define{{.*}} void @func_name +// LLVM-SPV: call{{.*}} void @take(ptr addrspace(2) noundef @__func__.func_name) + +// OGCG-SPV-LABEL: define{{.*}} void @func_name +// OGCG-SPV: call{{.*}} void @take(ptr addrspace(2) noundef @__func__.func_name) + +// LLVM-GCN-LABEL: define{{.*}} void @func_name +// LLVM-GCN: call void @take(ptr addrspace(4) noundef @__func__.func_name) + +// OGCG-GCN-LABEL: define{{.*}} void @func_name +// OGCG-GCN: call void @take(ptr addrspace(4) noundef @__func__.func_name) +void func_name(void) { take(__func__); } diff --git a/clang/test/CIR/CodeGenSYCL/string-literal-addrspace-nyi.cpp b/clang/test/CIR/CodeGenSYCL/string-literal-addrspace-nyi.cpp new file mode 100644 index 0000000000000..c56a1aac5dd78 --- /dev/null +++ b/clang/test/CIR/CodeGenSYCL/string-literal-addrspace-nyi.cpp @@ -0,0 +1,20 @@ +// RUN: %clang_cc1 -std=c++20 -fsycl-is-device -triple spirv64-unknown-unknown -fclangir -emit-cir %s -o /dev/null -verify + +// SYCL device string literals belong in sycl_global, which CIR cannot +// represent yet. + +// Required by sycl_kernel_entry_point semantics. +template <typename KernelName, typename... Ts> +void sycl_kernel_launch(const char *, Ts...) {} + +template <typename KernelName, typename KernelType> +[[clang::sycl_kernel_entry_point(KernelName)]] +void kernel_single_task(KernelType kf) { kf(); } + +struct KN; +void use(const char *); + +void test() { + // expected-error@*:* {{ClangIR code gen Not Yet Implemented: SYCL global constant address space}} + kernel_single_task<KN>([]() { use("hello"); }); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
