https://github.com/koparasy updated https://github.com/llvm/llvm-project/pull/229281
>From 93aa95805dde7ac9004d46a7cef6e47831aac6ee Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris <[email protected]> Date: Mon, 5 Oct 2026 17:35:17 -0700 Subject: [PATCH 1/4] [CIR][SYCL] Cast stack allocas to the generic address space Resolve the default address space of stack allocas through getTargetAddressSpace, as classic CodeGen does, so SYCL device locals on SPIR are accessed through the generic address space instead of leaking private pointers into generic ones. Also store parameters through the cast address and keep the destination address space in emitArrayInit. Co-Authored-By: Claude Opus 5.5 (1M context) <[email protected]> --- clang/lib/CIR/CodeGen/CIRGenExpr.cpp | 9 + clang/lib/CIR/CodeGen/CIRGenExprAggregate.cpp | 6 +- clang/lib/CIR/CodeGen/CIRGenFunction.cpp | 3 +- .../CIR/CodeGenHIP/amdgcnspirv-builtins.hip | 6 +- .../CodeGenSYCL/address-space-conversions.cpp | 408 ++++++++++++++++++ .../CodeGenSYCL/address-space-deduction.cpp | 171 ++++++++ .../CIR/CodeGenSYCL/local-temporaries.cpp | 229 ++++++++++ .../CIR/CodeGenSYCL/record-address-space.cpp | 2 +- 8 files changed, 825 insertions(+), 9 deletions(-) create mode 100644 clang/test/CIR/CodeGenSYCL/address-space-conversions.cpp create mode 100644 clang/test/CIR/CodeGenSYCL/address-space-deduction.cpp create mode 100644 clang/test/CIR/CodeGenSYCL/local-temporaries.cpp diff --git a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp index 51533f9dc0f9c..f1dcfe4a08bec 100644 --- a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp @@ -2952,6 +2952,15 @@ Address CIRGenFunction::maybeCastStackAddressSpace( destAddrSpace = cir::toCIRAddressSpaceAttr( getMLIRContext(), cgm.getLangTempAllocaAddressSpace()); + // Resolve the default address space through getTargetAddressSpace, as + // classic CodeGen does and as CIRGenTypes::getPointerAddressSpace does for + // default pointer types. This is only non-zero for targets where the default + // address space is not 0 (e.g. generic for SYCL device code). + if (!cir::normalizeDefaultAddressSpace(destAddrSpace)) + if (unsigned targetAS = getContext().getTargetAddressSpace(LangAS::Default)) + destAddrSpace = + cir::TargetAddressSpaceAttr::get(&getMLIRContext(), targetAS); + mlir::ptr::MemorySpaceAttrInterface srcAddrSpace = getCIRAllocaAddressSpace(); // Alloca always returns a pointer in alloca address space, which may // be different from the type defined by the language. For example, diff --git a/clang/lib/CIR/CodeGen/CIRGenExprAggregate.cpp b/clang/lib/CIR/CodeGen/CIRGenExprAggregate.cpp index 43f7842b03b3c..9e091568d2a7f 100644 --- a/clang/lib/CIR/CodeGen/CIRGenExprAggregate.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenExprAggregate.cpp @@ -751,8 +751,10 @@ void AggExprEmitter::emitArrayInit(Address destPtr, cir::ArrayType arrayTy, const QualType elementPtrType = cgf.getContext().getPointerType(elementType); const mlir::Type cirElementType = cgf.convertType(elementType); - const cir::PointerType cirElementPtrType = - builder.getPointerTo(cirElementType); + const cir::PointerType cirElementPtrType = builder.getPointerTo( + cirElementType, + mlir::cast<cir::PointerType>(destPtr.getPointer().getType()) + .getAddrSpace()); auto begin = cir::CastOp::create(builder, loc, cirElementPtrType, cir::CastKind::array_to_ptrdecay, diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.cpp b/clang/lib/CIR/CodeGen/CIRGenFunction.cpp index e5e83a2647018..0c7d014bbb6b7 100644 --- a/clang/lib/CIR/CodeGen/CIRGenFunction.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenFunction.cpp @@ -496,7 +496,8 @@ void CIRGenFunction::emitFunctionProlog(const FunctionArgList &args, // Location of the store to the param storage tracked as beginning of // the function body. mlir::Location fnBodyBegin = getLoc(bodyBeginLoc); - builder.CIRBaseBuilderTy::createStore(fnBodyBegin, paramVal, addrVal); + builder.CIRBaseBuilderTy::createStore(fnBodyBegin, paramVal, + addr.getPointer()); } assert(builder.getInsertionBlock() && "Should be valid"); } diff --git a/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip b/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip index dd92cfab285f1..eced98bad1210 100644 --- a/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip +++ b/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip @@ -16,10 +16,6 @@ // FIXME: CIR doesn't propagate the 'contract' fast-math flag to LLVM IR yet. // The LLVM checks match the flag-free form and fail once it lands. -// FIXME: CIR lowers the Default address space to 0 rather than generic on -// SPIR-V, so the LLVM output lacks the addrspacecast that classic codegen -// emits. The LLVM checks fail once it lands. - #define __device__ __attribute__((device)) __device__ int test_readfirstlane(int x) { @@ -30,7 +26,7 @@ __device__ int test_readfirstlane(int x) { // CIR: cir.call_llvm_intrinsic "amdgcn.readfirstlane" {{.*}} : (!s32i) -> !s32i // LLVM-LABEL: define spir_func noundef i32 @_Z18test_readfirstlanei -// LLVM-NOT: addrspacecast +// LLVM: addrspacecast ptr %{{.*}} to ptr addrspace(4) // LLVM: call{{.*}} @llvm.amdgcn.readfirstlane.i32 // OGCG-LABEL: define spir_func noundef i32 @_Z18test_readfirstlanei diff --git a/clang/test/CIR/CodeGenSYCL/address-space-conversions.cpp b/clang/test/CIR/CodeGenSYCL/address-space-conversions.cpp new file mode 100644 index 0000000000000..4ea1e603cd87a --- /dev/null +++ b/clang/test/CIR/CodeGenSYCL/address-space-conversions.cpp @@ -0,0 +1,408 @@ +// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -disable-llvm-passes -Wno-deprecated-attributes -fclangir -emit-cir %s -o %t.cir +// RUN: FileCheck --input-file=%t.cir %s -check-prefix=CIR +// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -disable-llvm-passes -Wno-deprecated-attributes -fclangir -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --input-file=%t-cir.ll %s -check-prefix=LLVM +// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -disable-llvm-passes -Wno-deprecated-attributes -emit-llvm %s -o %t.ll +// RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG + +// Port of clang/test/CodeGenSYCL/address-space-conversions.cpp. Local +// variables and parameters are allocated in the private address space and +// accessed through the generic address space. +// +// TODO(cir): Calls are missing the spir_func calling convention and null +// pointers to named address spaces are not emitted as an addrspacecast of the +// generic null pointer. + +void bar(int &Data) {} +void bar2(int &Data) {} +void bar(int [[clang::sycl_local]] &Data) {} +void foo(int *Data) {} +void foo2(int *Data) {} +void foo(int [[clang::sycl_local]] *Data) {} + +template <typename T> +void tmpl(T t) {} + +[[clang::sycl_external]] void usages() { + int *NoAS; + int [[clang::sycl_global]] *GLOB; + int [[clang::sycl_local]] *LOC; + int [[clang::sycl_private]] *PRIV; + int __attribute__((opencl_global_device)) *GLOBDEVICE; + int __attribute__((opencl_global_host)) *GLOBHOST; + + LOC = nullptr; + GLOB = nullptr; + + // Explicit conversions + // From named address spaces to default address space + NoAS = (int *)GLOB; + NoAS = (int *)LOC; + NoAS = (int *)PRIV; + // From default address space to named address space + GLOB = (int [[clang::sycl_global]] *)NoAS; + LOC = (int [[clang::sycl_local]] *)NoAS; + PRIV = (int [[clang::sycl_private]] *)NoAS; + // From opencl_global_[host/device] address spaces to sycl_global + GLOB = (int [[clang::sycl_global]] *)GLOBDEVICE; + GLOB = (int [[clang::sycl_global]] *)GLOBHOST; + + bar(*GLOB); + bar2(*GLOB); + + bar(*LOC); + bar2(*LOC); + + bar(*NoAS); + bar2(*NoAS); + + foo(GLOB); + foo2(GLOB); + foo(LOC); + foo2(LOC); + foo(NoAS); + foo2(NoAS); + + // Ensure that we still get 3 different template instantiations. + tmpl(GLOB); + tmpl(LOC); + tmpl(PRIV); + tmpl(NoAS); +} + + +// CIR-LABEL: cir.func {{.*}}@_Z6usagesv( +// CIR-NEXT: %[[NOAS:.*]] = cir.alloca "NoAS" align(8) : !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>> +// CIR-NEXT: %[[GLOB:.*]] = cir.alloca "GLOB" align(8) : !cir.ptr<!cir.ptr<!s32i, target_address_space(1)>> +// CIR-NEXT: %[[LOC:.*]] = cir.alloca "LOC" align(8) : !cir.ptr<!cir.ptr<!s32i, target_address_space(3)>> +// CIR-NEXT: %[[PRIV:.*]] = cir.alloca "PRIV" align(8) : !cir.ptr<!cir.ptr<!s32i>> +// CIR-NEXT: %[[GLOBDEVICE:.*]] = cir.alloca "GLOBDEVICE" align(8) : !cir.ptr<!cir.ptr<!s32i, target_address_space(5)>> +// CIR-NEXT: %[[GLOBHOST:.*]] = cir.alloca "GLOBHOST" align(8) : !cir.ptr<!cir.ptr<!s32i, target_address_space(6)>> +// CIR-NEXT: %[[GLOBHOST_ASCAST:.*]] = cir.cast address_space %[[GLOBHOST]] : !cir.ptr<!cir.ptr<!s32i, target_address_space(6)>> -> !cir.ptr<!cir.ptr<!s32i, target_address_space(6)>, target_address_space(4)> +// CIR-NEXT: %[[GLOBDEVICE_ASCAST:.*]] = cir.cast address_space %[[GLOBDEVICE]] : !cir.ptr<!cir.ptr<!s32i, target_address_space(5)>> -> !cir.ptr<!cir.ptr<!s32i, target_address_space(5)>, target_address_space(4)> +// CIR-NEXT: %[[PRIV_ASCAST:.*]] = cir.cast address_space %[[PRIV]] : !cir.ptr<!cir.ptr<!s32i>> -> !cir.ptr<!cir.ptr<!s32i>, target_address_space(4)> +// CIR-NEXT: %[[LOC_ASCAST:.*]] = cir.cast address_space %[[LOC]] : !cir.ptr<!cir.ptr<!s32i, target_address_space(3)>> -> !cir.ptr<!cir.ptr<!s32i, target_address_space(3)>, target_address_space(4)> +// CIR-NEXT: %[[GLOB_ASCAST:.*]] = cir.cast address_space %[[GLOB]] : !cir.ptr<!cir.ptr<!s32i, target_address_space(1)>> -> !cir.ptr<!cir.ptr<!s32i, target_address_space(1)>, target_address_space(4)> +// CIR-NEXT: %[[NOAS_ASCAST:.*]] = cir.cast address_space %[[NOAS]] : !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>> -> !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>, target_address_space(4)> + +// CIR-LABEL: cir.func {{.*}}@_Z3barRi( +// CIR-NEXT: %[[DATA:.*]] = cir.alloca "Data" align(8) init const : !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>> +// CIR-NEXT: %[[DATA_ASCAST:.*]] = cir.cast address_space %[[DATA]] : !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>> -> !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>, target_address_space(4)> +// CIR-NEXT: cir.store %arg0, %[[DATA_ASCAST]] : !cir.ptr<!s32i, target_address_space(4)>, !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>, target_address_space(4)> +// CIR-NEXT: cir.return + +// CIR-LABEL: cir.func {{.*}}@_Z3barRU3AS3i( +// CIR-NEXT: %[[DATA:.*]] = cir.alloca "Data" align(8) init const : !cir.ptr<!cir.ptr<!s32i, target_address_space(3)>> +// CIR-NEXT: %[[DATA_ASCAST:.*]] = cir.cast address_space %[[DATA]] : !cir.ptr<!cir.ptr<!s32i, target_address_space(3)>> -> !cir.ptr<!cir.ptr<!s32i, target_address_space(3)>, target_address_space(4)> +// CIR-NEXT: cir.store %arg0, %[[DATA_ASCAST]] : !cir.ptr<!s32i, target_address_space(3)>, !cir.ptr<!cir.ptr<!s32i, target_address_space(3)>, target_address_space(4)> +// CIR-NEXT: cir.return + +// CIR-LABEL: cir.func {{.*}}@_Z3fooPi( +// CIR-NEXT: %[[DATA:.*]] = cir.alloca "Data" align(8) init : !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>> +// CIR-NEXT: %[[DATA_ASCAST:.*]] = cir.cast address_space %[[DATA]] : !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>> -> !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>, target_address_space(4)> +// CIR-NEXT: cir.store %arg0, %[[DATA_ASCAST]] : !cir.ptr<!s32i, target_address_space(4)>, !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>, target_address_space(4)> +// CIR-NEXT: cir.return + +// CIR-LABEL: cir.func {{.*}}@_Z3fooPU3AS3i( +// CIR-NEXT: %[[DATA:.*]] = cir.alloca "Data" align(8) init : !cir.ptr<!cir.ptr<!s32i, target_address_space(3)>> +// CIR-NEXT: %[[DATA_ASCAST:.*]] = cir.cast address_space %[[DATA]] : !cir.ptr<!cir.ptr<!s32i, target_address_space(3)>> -> !cir.ptr<!cir.ptr<!s32i, target_address_space(3)>, target_address_space(4)> +// CIR-NEXT: cir.store %arg0, %[[DATA_ASCAST]] : !cir.ptr<!s32i, target_address_space(3)>, !cir.ptr<!cir.ptr<!s32i, target_address_space(3)>, target_address_space(4)> +// CIR-NEXT: cir.return + +// LLVM-LABEL: define {{.*}}spir_func void @_Z6usagesv( +// LLVM: %[[NOAS:.*]] = alloca ptr addrspace(4), align 8 +// LLVM-NEXT: %[[GLOB:.*]] = alloca ptr addrspace(1), align 8 +// LLVM-NEXT: %[[LOC:.*]] = alloca ptr addrspace(3), align 8 +// LLVM-NEXT: %[[PRIV:.*]] = alloca ptr, align 8 +// LLVM-NEXT: %[[GLOBDEVICE:.*]] = alloca ptr addrspace(5), align 8 +// LLVM-NEXT: %[[GLOBHOST:.*]] = alloca ptr addrspace(6), align 8 +// LLVM-NEXT: %[[GLOBHOST_ASCAST:.*]] = addrspacecast ptr %[[GLOBHOST]] to ptr addrspace(4) +// LLVM-NEXT: %[[GLOBDEVICE_ASCAST:.*]] = addrspacecast ptr %[[GLOBDEVICE]] to ptr addrspace(4) +// LLVM-NEXT: %[[PRIV_ASCAST:.*]] = addrspacecast ptr %[[PRIV]] to ptr addrspace(4) +// LLVM-NEXT: %[[LOC_ASCAST:.*]] = addrspacecast ptr %[[LOC]] to ptr addrspace(4) +// LLVM-NEXT: %[[GLOB_ASCAST:.*]] = addrspacecast ptr %[[GLOB]] to ptr addrspace(4) +// LLVM-NEXT: %[[NOAS_ASCAST:.*]] = addrspacecast ptr %[[NOAS]] to ptr addrspace(4) +// LLVM-NEXT: store ptr addrspace(3) null, ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// LLVM-NEXT: store ptr addrspace(1) null, ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP0:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP1:.*]] = addrspacecast ptr addrspace(1) %[[TMP0]] to ptr addrspace(4) +// LLVM-NEXT: store ptr addrspace(4) %[[TMP1]], ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP2:.*]] = load ptr addrspace(3), ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP3:.*]] = addrspacecast ptr addrspace(3) %[[TMP2]] to ptr addrspace(4) +// LLVM-NEXT: store ptr addrspace(4) %[[TMP3]], ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP4:.*]] = load ptr, ptr addrspace(4) %[[PRIV_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP5:.*]] = addrspacecast ptr %[[TMP4]] to ptr addrspace(4) +// LLVM-NEXT: store ptr addrspace(4) %[[TMP5]], ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP6:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP7:.*]] = addrspacecast ptr addrspace(4) %[[TMP6]] to ptr addrspace(1) +// LLVM-NEXT: store ptr addrspace(1) %[[TMP7]], ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP8:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP9:.*]] = addrspacecast ptr addrspace(4) %[[TMP8]] to ptr addrspace(3) +// LLVM-NEXT: store ptr addrspace(3) %[[TMP9]], ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP10:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP11:.*]] = addrspacecast ptr addrspace(4) %[[TMP10]] to ptr +// LLVM-NEXT: store ptr %[[TMP11]], ptr addrspace(4) %[[PRIV_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP12:.*]] = load ptr addrspace(5), ptr addrspace(4) %[[GLOBDEVICE_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP13:.*]] = addrspacecast ptr addrspace(5) %[[TMP12]] to ptr addrspace(1) +// LLVM-NEXT: store ptr addrspace(1) %[[TMP13]], ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP14:.*]] = load ptr addrspace(6), ptr addrspace(4) %[[GLOBHOST_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP15:.*]] = addrspacecast ptr addrspace(6) %[[TMP14]] to ptr addrspace(1) +// LLVM-NEXT: store ptr addrspace(1) %[[TMP15]], ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP16:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP17:.*]] = addrspacecast ptr addrspace(1) %[[TMP16]] to ptr addrspace(4) +// LLVM-NEXT: call void @_Z3barRi(ptr addrspace(4) noundef align 4 dereferenceable(4) %[[TMP17]]) +// LLVM-NEXT: %[[TMP18:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP19:.*]] = addrspacecast ptr addrspace(1) %[[TMP18]] to ptr addrspace(4) +// LLVM-NEXT: call void @_Z4bar2Ri(ptr addrspace(4) noundef align 4 dereferenceable(4) %[[TMP19]]) +// LLVM-NEXT: %[[TMP20:.*]] = load ptr addrspace(3), ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// LLVM-NEXT: call void @_Z3barRU3AS3i(ptr addrspace(3) noundef align 4 dereferenceable(4) %[[TMP20]]) +// LLVM-NEXT: %[[TMP21:.*]] = load ptr addrspace(3), ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP22:.*]] = addrspacecast ptr addrspace(3) %[[TMP21]] to ptr addrspace(4) +// LLVM-NEXT: call void @_Z4bar2Ri(ptr addrspace(4) noundef align 4 dereferenceable(4) %[[TMP22]]) +// LLVM-NEXT: %[[TMP23:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// LLVM-NEXT: call void @_Z3barRi(ptr addrspace(4) noundef align 4 dereferenceable(4) %[[TMP23]]) +// LLVM-NEXT: %[[TMP24:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// LLVM-NEXT: call void @_Z4bar2Ri(ptr addrspace(4) noundef align 4 dereferenceable(4) %[[TMP24]]) +// LLVM-NEXT: %[[TMP25:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP26:.*]] = addrspacecast ptr addrspace(1) %[[TMP25]] to ptr addrspace(4) +// LLVM-NEXT: call void @_Z3fooPi(ptr addrspace(4) noundef %[[TMP26]]) +// LLVM-NEXT: %[[TMP27:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP28:.*]] = addrspacecast ptr addrspace(1) %[[TMP27]] to ptr addrspace(4) +// LLVM-NEXT: call void @_Z4foo2Pi(ptr addrspace(4) noundef %[[TMP28]]) +// LLVM-NEXT: %[[TMP29:.*]] = load ptr addrspace(3), ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// LLVM-NEXT: call void @_Z3fooPU3AS3i(ptr addrspace(3) noundef %[[TMP29]]) +// LLVM-NEXT: %[[TMP30:.*]] = load ptr addrspace(3), ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP31:.*]] = addrspacecast ptr addrspace(3) %[[TMP30]] to ptr addrspace(4) +// LLVM-NEXT: call void @_Z4foo2Pi(ptr addrspace(4) noundef %[[TMP31]]) +// LLVM-NEXT: %[[TMP32:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// LLVM-NEXT: call void @_Z3fooPi(ptr addrspace(4) noundef %[[TMP32]]) +// LLVM-NEXT: %[[TMP33:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// LLVM-NEXT: call void @_Z4foo2Pi(ptr addrspace(4) noundef %[[TMP33]]) +// LLVM-NEXT: %[[TMP34:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// LLVM-NEXT: call void @_Z4tmplIPU3AS1iEvT_(ptr addrspace(1) noundef %[[TMP34]]) +// LLVM-NEXT: %[[TMP35:.*]] = load ptr addrspace(3), ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// LLVM-NEXT: call void @_Z4tmplIPU3AS3iEvT_(ptr addrspace(3) noundef %[[TMP35]]) +// LLVM-NEXT: %[[TMP36:.*]] = load ptr, ptr addrspace(4) %[[PRIV_ASCAST]], align 8 +// LLVM-NEXT: call void @_Z4tmplIPU3AS0iEvT_(ptr noundef %[[TMP36]]) +// LLVM-NEXT: %[[TMP37:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// LLVM-NEXT: call void @_Z4tmplIPiEvT_(ptr addrspace(4) noundef %[[TMP37]]) +// LLVM-NEXT: ret void + +// LLVM-LABEL: define {{.*}}spir_func void @_Z3barRi( +// LLVM-SAME: ptr addrspace(4) noundef align 4 dereferenceable(4) %[[DATA:.*]]) +// LLVM: %[[DATA_ADDR:.*]] = alloca ptr addrspace(4), align 8 +// LLVM-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) +// LLVM-NEXT: store ptr addrspace(4) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 +// LLVM-NEXT: ret void + +// LLVM-LABEL: define {{.*}}spir_func void @_Z4bar2Ri( +// LLVM-SAME: ptr addrspace(4) noundef align 4 dereferenceable(4) %[[DATA:.*]]) +// LLVM: %[[DATA_ADDR:.*]] = alloca ptr addrspace(4), align 8 +// LLVM-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) +// LLVM-NEXT: store ptr addrspace(4) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 +// LLVM-NEXT: ret void + +// LLVM-LABEL: define {{.*}}spir_func void @_Z3barRU3AS3i( +// LLVM-SAME: ptr addrspace(3) noundef align 4 dereferenceable(4) %[[DATA:.*]]) +// LLVM: %[[DATA_ADDR:.*]] = alloca ptr addrspace(3), align 8 +// LLVM-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) +// LLVM-NEXT: store ptr addrspace(3) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 +// LLVM-NEXT: ret void + +// LLVM-LABEL: define {{.*}}spir_func void @_Z3fooPi( +// LLVM-SAME: ptr addrspace(4) noundef %[[DATA:.*]]) +// LLVM: %[[DATA_ADDR:.*]] = alloca ptr addrspace(4), align 8 +// LLVM-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) +// LLVM-NEXT: store ptr addrspace(4) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 +// LLVM-NEXT: ret void + +// LLVM-LABEL: define {{.*}}spir_func void @_Z4foo2Pi( +// LLVM-SAME: ptr addrspace(4) noundef %[[DATA:.*]]) +// LLVM: %[[DATA_ADDR:.*]] = alloca ptr addrspace(4), align 8 +// LLVM-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) +// LLVM-NEXT: store ptr addrspace(4) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 +// LLVM-NEXT: ret void + +// LLVM-LABEL: define {{.*}}spir_func void @_Z3fooPU3AS3i( +// LLVM-SAME: ptr addrspace(3) noundef %[[DATA:.*]]) +// LLVM: %[[DATA_ADDR:.*]] = alloca ptr addrspace(3), align 8 +// LLVM-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) +// LLVM-NEXT: store ptr addrspace(3) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 +// LLVM-NEXT: ret void + +// LLVM-LABEL: define {{.*}}spir_func void @_Z4tmplIPU3AS1iEvT_( +// LLVM-SAME: ptr addrspace(1) noundef %[[T:.*]]) +// LLVM: %[[T_ADDR:.*]] = alloca ptr addrspace(1), align 8 +// LLVM-NEXT: %[[T_ADDR_ASCAST:.*]] = addrspacecast ptr %[[T_ADDR]] to ptr addrspace(4) +// LLVM-NEXT: store ptr addrspace(1) %[[T]], ptr addrspace(4) %[[T_ADDR_ASCAST]], align 8 +// LLVM-NEXT: ret void + +// LLVM-LABEL: define {{.*}}spir_func void @_Z4tmplIPU3AS3iEvT_( +// LLVM-SAME: ptr addrspace(3) noundef %[[T:.*]]) +// LLVM: %[[T_ADDR:.*]] = alloca ptr addrspace(3), align 8 +// LLVM-NEXT: %[[T_ADDR_ASCAST:.*]] = addrspacecast ptr %[[T_ADDR]] to ptr addrspace(4) +// LLVM-NEXT: store ptr addrspace(3) %[[T]], ptr addrspace(4) %[[T_ADDR_ASCAST]], align 8 +// LLVM-NEXT: ret void + +// LLVM-LABEL: define {{.*}}spir_func void @_Z4tmplIPU3AS0iEvT_( +// LLVM-SAME: ptr noundef %[[T:.*]]) +// LLVM: %[[T_ADDR:.*]] = alloca ptr, align 8 +// LLVM-NEXT: %[[T_ADDR_ASCAST:.*]] = addrspacecast ptr %[[T_ADDR]] to ptr addrspace(4) +// LLVM-NEXT: store ptr %[[T]], ptr addrspace(4) %[[T_ADDR_ASCAST]], align 8 +// LLVM-NEXT: ret void + +// LLVM-LABEL: define {{.*}}spir_func void @_Z4tmplIPiEvT_( +// LLVM-SAME: ptr addrspace(4) noundef %[[T:.*]]) +// LLVM: %[[T_ADDR:.*]] = alloca ptr addrspace(4), align 8 +// LLVM-NEXT: %[[T_ADDR_ASCAST:.*]] = addrspacecast ptr %[[T_ADDR]] to ptr addrspace(4) +// LLVM-NEXT: store ptr addrspace(4) %[[T]], ptr addrspace(4) %[[T_ADDR_ASCAST]], align 8 +// LLVM-NEXT: ret void + +// OGCG-LABEL: define {{.*}}spir_func void @_Z6usagesv( +// OGCG: %[[NOAS:.*]] = alloca ptr addrspace(4), align 8 +// OGCG-NEXT: %[[GLOB:.*]] = alloca ptr addrspace(1), align 8 +// OGCG-NEXT: %[[LOC:.*]] = alloca ptr addrspace(3), align 8 +// OGCG-NEXT: %[[PRIV:.*]] = alloca ptr, align 8 +// OGCG-NEXT: %[[GLOBDEVICE:.*]] = alloca ptr addrspace(5), align 8 +// OGCG-NEXT: %[[GLOBHOST:.*]] = alloca ptr addrspace(6), align 8 +// OGCG-NEXT: %[[NOAS_ASCAST:.*]] = addrspacecast ptr %[[NOAS]] to ptr addrspace(4) +// OGCG-NEXT: %[[GLOB_ASCAST:.*]] = addrspacecast ptr %[[GLOB]] to ptr addrspace(4) +// OGCG-NEXT: %[[LOC_ASCAST:.*]] = addrspacecast ptr %[[LOC]] to ptr addrspace(4) +// OGCG-NEXT: %[[PRIV_ASCAST:.*]] = addrspacecast ptr %[[PRIV]] to ptr addrspace(4) +// OGCG-NEXT: %[[GLOBDEVICE_ASCAST:.*]] = addrspacecast ptr %[[GLOBDEVICE]] to ptr addrspace(4) +// OGCG-NEXT: %[[GLOBHOST_ASCAST:.*]] = addrspacecast ptr %[[GLOBHOST]] to ptr addrspace(4) +// OGCG-NEXT: store ptr addrspace(3) addrspacecast (ptr addrspace(4) null to ptr addrspace(3)), ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// OGCG-NEXT: store ptr addrspace(1) addrspacecast (ptr addrspace(4) null to ptr addrspace(1)), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP0:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP1:.*]] = addrspacecast ptr addrspace(1) %[[TMP0]] to ptr addrspace(4) +// OGCG-NEXT: store ptr addrspace(4) %[[TMP1]], ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP2:.*]] = load ptr addrspace(3), ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP3:.*]] = addrspacecast ptr addrspace(3) %[[TMP2]] to ptr addrspace(4) +// OGCG-NEXT: store ptr addrspace(4) %[[TMP3]], ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP4:.*]] = load ptr, ptr addrspace(4) %[[PRIV_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP5:.*]] = addrspacecast ptr %[[TMP4]] to ptr addrspace(4) +// OGCG-NEXT: store ptr addrspace(4) %[[TMP5]], ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP6:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP7:.*]] = addrspacecast ptr addrspace(4) %[[TMP6]] to ptr addrspace(1) +// OGCG-NEXT: store ptr addrspace(1) %[[TMP7]], ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP8:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP9:.*]] = addrspacecast ptr addrspace(4) %[[TMP8]] to ptr addrspace(3) +// OGCG-NEXT: store ptr addrspace(3) %[[TMP9]], ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP10:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP11:.*]] = addrspacecast ptr addrspace(4) %[[TMP10]] to ptr +// OGCG-NEXT: store ptr %[[TMP11]], ptr addrspace(4) %[[PRIV_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP12:.*]] = load ptr addrspace(5), ptr addrspace(4) %[[GLOBDEVICE_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP13:.*]] = addrspacecast ptr addrspace(5) %[[TMP12]] to ptr addrspace(1) +// OGCG-NEXT: store ptr addrspace(1) %[[TMP13]], ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP14:.*]] = load ptr addrspace(6), ptr addrspace(4) %[[GLOBHOST_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP15:.*]] = addrspacecast ptr addrspace(6) %[[TMP14]] to ptr addrspace(1) +// OGCG-NEXT: store ptr addrspace(1) %[[TMP15]], ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP16:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP17:.*]] = addrspacecast ptr addrspace(1) %[[TMP16]] to ptr addrspace(4) +// OGCG-NEXT: call spir_func void @_Z3barRi(ptr addrspace(4) noundef align 4 dereferenceable(4) %[[TMP17]]) +// OGCG-NEXT: %[[TMP18:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP19:.*]] = addrspacecast ptr addrspace(1) %[[TMP18]] to ptr addrspace(4) +// OGCG-NEXT: call spir_func void @_Z4bar2Ri(ptr addrspace(4) noundef align 4 dereferenceable(4) %[[TMP19]]) +// OGCG-NEXT: %[[TMP20:.*]] = load ptr addrspace(3), ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// OGCG-NEXT: call spir_func void @_Z3barRU3AS3i(ptr addrspace(3) noundef align 4 dereferenceable(4) %[[TMP20]]) +// OGCG-NEXT: %[[TMP21:.*]] = load ptr addrspace(3), ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP22:.*]] = addrspacecast ptr addrspace(3) %[[TMP21]] to ptr addrspace(4) +// OGCG-NEXT: call spir_func void @_Z4bar2Ri(ptr addrspace(4) noundef align 4 dereferenceable(4) %[[TMP22]]) +// OGCG-NEXT: %[[TMP23:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// OGCG-NEXT: call spir_func void @_Z3barRi(ptr addrspace(4) noundef align 4 dereferenceable(4) %[[TMP23]]) +// OGCG-NEXT: %[[TMP24:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// OGCG-NEXT: call spir_func void @_Z4bar2Ri(ptr addrspace(4) noundef align 4 dereferenceable(4) %[[TMP24]]) +// OGCG-NEXT: %[[TMP25:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP26:.*]] = addrspacecast ptr addrspace(1) %[[TMP25]] to ptr addrspace(4) +// OGCG-NEXT: call spir_func void @_Z3fooPi(ptr addrspace(4) noundef %[[TMP26]]) +// OGCG-NEXT: %[[TMP27:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP28:.*]] = addrspacecast ptr addrspace(1) %[[TMP27]] to ptr addrspace(4) +// OGCG-NEXT: call spir_func void @_Z4foo2Pi(ptr addrspace(4) noundef %[[TMP28]]) +// OGCG-NEXT: %[[TMP29:.*]] = load ptr addrspace(3), ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// OGCG-NEXT: call spir_func void @_Z3fooPU3AS3i(ptr addrspace(3) noundef %[[TMP29]]) +// OGCG-NEXT: %[[TMP30:.*]] = load ptr addrspace(3), ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP31:.*]] = addrspacecast ptr addrspace(3) %[[TMP30]] to ptr addrspace(4) +// OGCG-NEXT: call spir_func void @_Z4foo2Pi(ptr addrspace(4) noundef %[[TMP31]]) +// OGCG-NEXT: %[[TMP32:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// OGCG-NEXT: call spir_func void @_Z3fooPi(ptr addrspace(4) noundef %[[TMP32]]) +// OGCG-NEXT: %[[TMP33:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// OGCG-NEXT: call spir_func void @_Z4foo2Pi(ptr addrspace(4) noundef %[[TMP33]]) +// OGCG-NEXT: %[[TMP34:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// OGCG-NEXT: call spir_func void @_Z4tmplIPU3AS1iEvT_(ptr addrspace(1) noundef %[[TMP34]]) +// OGCG-NEXT: %[[TMP35:.*]] = load ptr addrspace(3), ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// OGCG-NEXT: call spir_func void @_Z4tmplIPU3AS3iEvT_(ptr addrspace(3) noundef %[[TMP35]]) +// OGCG-NEXT: %[[TMP36:.*]] = load ptr, ptr addrspace(4) %[[PRIV_ASCAST]], align 8 +// OGCG-NEXT: call spir_func void @_Z4tmplIPU3AS0iEvT_(ptr noundef %[[TMP36]]) +// OGCG-NEXT: %[[TMP37:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// OGCG-NEXT: call spir_func void @_Z4tmplIPiEvT_(ptr addrspace(4) noundef %[[TMP37]]) +// OGCG-NEXT: ret void + +// OGCG-LABEL: define {{.*}}spir_func void @_Z3barRi( +// OGCG-SAME: ptr addrspace(4) noundef align 4 dereferenceable(4) %[[DATA:.*]]) +// OGCG: %[[DATA_ADDR:.*]] = alloca ptr addrspace(4), align 8 +// OGCG-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) +// OGCG-NEXT: store ptr addrspace(4) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 +// OGCG-NEXT: ret void + +// OGCG-LABEL: define {{.*}}spir_func void @_Z4bar2Ri( +// OGCG-SAME: ptr addrspace(4) noundef align 4 dereferenceable(4) %[[DATA:.*]]) +// OGCG: %[[DATA_ADDR:.*]] = alloca ptr addrspace(4), align 8 +// OGCG-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) +// OGCG-NEXT: store ptr addrspace(4) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 +// OGCG-NEXT: ret void + +// OGCG-LABEL: define {{.*}}spir_func void @_Z3barRU3AS3i( +// OGCG-SAME: ptr addrspace(3) noundef align 4 dereferenceable(4) %[[DATA:.*]]) +// OGCG: %[[DATA_ADDR:.*]] = alloca ptr addrspace(3), align 8 +// OGCG-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) +// OGCG-NEXT: store ptr addrspace(3) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 +// OGCG-NEXT: ret void + +// OGCG-LABEL: define {{.*}}spir_func void @_Z3fooPi( +// OGCG-SAME: ptr addrspace(4) noundef %[[DATA:.*]]) +// OGCG: %[[DATA_ADDR:.*]] = alloca ptr addrspace(4), align 8 +// OGCG-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) +// OGCG-NEXT: store ptr addrspace(4) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 +// OGCG-NEXT: ret void + +// OGCG-LABEL: define {{.*}}spir_func void @_Z4foo2Pi( +// OGCG-SAME: ptr addrspace(4) noundef %[[DATA:.*]]) +// OGCG: %[[DATA_ADDR:.*]] = alloca ptr addrspace(4), align 8 +// OGCG-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) +// OGCG-NEXT: store ptr addrspace(4) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 +// OGCG-NEXT: ret void + +// OGCG-LABEL: define {{.*}}spir_func void @_Z3fooPU3AS3i( +// OGCG-SAME: ptr addrspace(3) noundef %[[DATA:.*]]) +// OGCG: %[[DATA_ADDR:.*]] = alloca ptr addrspace(3), align 8 +// OGCG-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) +// OGCG-NEXT: store ptr addrspace(3) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 +// OGCG-NEXT: ret void + +// OGCG-LABEL: define {{.*}}spir_func void @_Z4tmplIPU3AS1iEvT_( +// OGCG-SAME: ptr addrspace(1) noundef %[[T:.*]]) +// OGCG: %[[T_ADDR:.*]] = alloca ptr addrspace(1), align 8 +// OGCG-NEXT: %[[T_ADDR_ASCAST:.*]] = addrspacecast ptr %[[T_ADDR]] to ptr addrspace(4) +// OGCG-NEXT: store ptr addrspace(1) %[[T]], ptr addrspace(4) %[[T_ADDR_ASCAST]], align 8 +// OGCG-NEXT: ret void + +// OGCG-LABEL: define {{.*}}spir_func void @_Z4tmplIPU3AS3iEvT_( +// OGCG-SAME: ptr addrspace(3) noundef %[[T:.*]]) +// OGCG: %[[T_ADDR:.*]] = alloca ptr addrspace(3), align 8 +// OGCG-NEXT: %[[T_ADDR_ASCAST:.*]] = addrspacecast ptr %[[T_ADDR]] to ptr addrspace(4) +// OGCG-NEXT: store ptr addrspace(3) %[[T]], ptr addrspace(4) %[[T_ADDR_ASCAST]], align 8 +// OGCG-NEXT: ret void + +// OGCG-LABEL: define {{.*}}spir_func void @_Z4tmplIPU3AS0iEvT_( +// OGCG-SAME: ptr noundef %[[T:.*]]) +// OGCG: %[[T_ADDR:.*]] = alloca ptr, align 8 +// OGCG-NEXT: %[[T_ADDR_ASCAST:.*]] = addrspacecast ptr %[[T_ADDR]] to ptr addrspace(4) +// OGCG-NEXT: store ptr %[[T]], ptr addrspace(4) %[[T_ADDR_ASCAST]], align 8 +// OGCG-NEXT: ret void + +// OGCG-LABEL: define {{.*}}spir_func void @_Z4tmplIPiEvT_( +// OGCG-SAME: ptr addrspace(4) noundef %[[T:.*]]) +// OGCG: %[[T_ADDR:.*]] = alloca ptr addrspace(4), align 8 +// OGCG-NEXT: %[[T_ADDR_ASCAST:.*]] = addrspacecast ptr %[[T_ADDR]] to ptr addrspace(4) +// OGCG-NEXT: store ptr addrspace(4) %[[T]], ptr addrspace(4) %[[T_ADDR_ASCAST]], align 8 +// OGCG-NEXT: ret void diff --git a/clang/test/CIR/CodeGenSYCL/address-space-deduction.cpp b/clang/test/CIR/CodeGenSYCL/address-space-deduction.cpp new file mode 100644 index 0000000000000..22032fcfa070e --- /dev/null +++ b/clang/test/CIR/CodeGenSYCL/address-space-deduction.cpp @@ -0,0 +1,171 @@ +// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -disable-llvm-passes -fclangir -emit-cir %s -o %t.cir +// RUN: FileCheck --input-file=%t.cir %s -check-prefix=CIR +// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -disable-llvm-passes -fclangir -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --input-file=%t-cir.ll %s -check-prefix=LLVM +// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o %t.ll +// RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG + +// Port of clang/test/CodeGenSYCL/address-space-deduction.cpp. Local variables +// are allocated in the private address space and cast to the generic address +// space, so taking their address yields a generic pointer. +// +// TODO(cir): The static local and the string literals of the original test +// are omitted until CIR supports the SYCL global address space. + +[[clang::sycl_external]] void test() { + int i = 0; + int *pptr = &i; + bool is_i_ptr = (pptr == &i); + + int var23 = 23; + char *cp = (char *)&var23; + *cp = 41; + + int arr[42]; + char *cpp = (char *)arr; + *cpp = 43; + + int *aptr = arr + 10; + if (aptr < arr + sizeof(arr)) + *aptr = 44; +} + +// CIR-LABEL: cir.func {{.*}}@_Z4testv() +// CIR: %[[I:.*]] = cir.alloca "i" {{.*}} : !cir.ptr<!s32i> +// CIR: %[[PPTR:.*]] = cir.alloca "pptr" {{.*}} : !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>> +// CIR: %[[IS_I_PTR:.*]] = cir.alloca "is_i_ptr" {{.*}} : !cir.ptr<!cir.bool> +// CIR: %[[VAR23:.*]] = cir.alloca "var23" {{.*}} : !cir.ptr<!s32i> +// CIR: %[[CP:.*]] = cir.alloca "cp" {{.*}} : !cir.ptr<!cir.ptr<!s8i, target_address_space(4)>> +// CIR: %[[ARR:.*]] = cir.alloca "arr" {{.*}} : !cir.ptr<!cir.array<!s32i x 42>> +// CIR: %[[CPP:.*]] = cir.alloca "cpp" {{.*}} : !cir.ptr<!cir.ptr<!s8i, target_address_space(4)>> +// CIR: %[[APTR:.*]] = cir.alloca "aptr" {{.*}} : !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>> +// CIR: %[[APTR_ASCAST:.*]] = cir.cast address_space %[[APTR]] : !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>> -> !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>, target_address_space(4)> +// CIR: %[[CPP_ASCAST:.*]] = cir.cast address_space %[[CPP]] : !cir.ptr<!cir.ptr<!s8i, target_address_space(4)>> -> !cir.ptr<!cir.ptr<!s8i, target_address_space(4)>, target_address_space(4)> +// CIR: %[[ARR_ASCAST:.*]] = cir.cast address_space %[[ARR]] : !cir.ptr<!cir.array<!s32i x 42>> -> !cir.ptr<!cir.array<!s32i x 42>, target_address_space(4)> +// CIR: %[[CP_ASCAST:.*]] = cir.cast address_space %[[CP]] : !cir.ptr<!cir.ptr<!s8i, target_address_space(4)>> -> !cir.ptr<!cir.ptr<!s8i, target_address_space(4)>, target_address_space(4)> +// CIR: %[[VAR23_ASCAST:.*]] = cir.cast address_space %[[VAR23]] : !cir.ptr<!s32i> -> !cir.ptr<!s32i, target_address_space(4)> +// CIR: %[[IS_I_PTR_ASCAST:.*]] = cir.cast address_space %[[IS_I_PTR]] : !cir.ptr<!cir.bool> -> !cir.ptr<!cir.bool, target_address_space(4)> +// CIR: %[[PPTR_ASCAST:.*]] = cir.cast address_space %[[PPTR]] : !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>> -> !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>, target_address_space(4)> +// CIR: %[[I_ASCAST:.*]] = cir.cast address_space %[[I]] : !cir.ptr<!s32i> -> !cir.ptr<!s32i, target_address_space(4)> +// CIR: cir.store {{.*}}, %[[I_ASCAST]] : !s32i, !cir.ptr<!s32i, target_address_space(4)> +// CIR: cir.store {{.*}} %[[I_ASCAST]], %[[PPTR_ASCAST]] : !cir.ptr<!s32i, target_address_space(4)>, !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>, target_address_space(4)> +// CIR: %[[PPTR_VAL:.*]] = cir.load {{.*}} %[[PPTR_ASCAST]] +// CIR: %[[CMP:.*]] = cir.cmp eq %[[PPTR_VAL]], %[[I_ASCAST]] : !cir.ptr<!s32i, target_address_space(4)> +// CIR: cir.store {{.*}} %[[CMP]], %[[IS_I_PTR_ASCAST]] : !cir.bool, !cir.ptr<!cir.bool, target_address_space(4)> +// CIR: cir.store {{.*}}, %[[VAR23_ASCAST]] : !s32i, !cir.ptr<!s32i, target_address_space(4)> +// CIR: %[[VAR23_CHAR:.*]] = cir.cast bitcast %[[VAR23_ASCAST]] : !cir.ptr<!s32i, target_address_space(4)> -> !cir.ptr<!s8i, target_address_space(4)> +// CIR: cir.store {{.*}} %[[VAR23_CHAR]], %[[CP_ASCAST]] +// CIR: %[[CP_VAL:.*]] = cir.load deref {{.*}} %[[CP_ASCAST]] +// CIR: cir.store {{.*}}, %[[CP_VAL]] : !s8i, !cir.ptr<!s8i, target_address_space(4)> +// CIR: %[[ARR_DECAY:.*]] = cir.cast array_to_ptrdecay %[[ARR_ASCAST]] : !cir.ptr<!cir.array<!s32i x 42>, target_address_space(4)> -> !cir.ptr<!s32i, target_address_space(4)> +// CIR: %[[ARR_CHAR:.*]] = cir.cast bitcast %[[ARR_DECAY]] : !cir.ptr<!s32i, target_address_space(4)> -> !cir.ptr<!s8i, target_address_space(4)> +// CIR: cir.store {{.*}} %[[ARR_CHAR]], %[[CPP_ASCAST]] +// CIR: %[[CPP_VAL:.*]] = cir.load deref {{.*}} %[[CPP_ASCAST]] +// CIR: cir.store {{.*}}, %[[CPP_VAL]] : !s8i, !cir.ptr<!s8i, target_address_space(4)> +// CIR: %[[ARR_DECAY1:.*]] = cir.cast array_to_ptrdecay %[[ARR_ASCAST]] +// CIR: %[[ADD_PTR:.*]] = cir.ptr_stride %[[ARR_DECAY1]], {{.*}} : (!cir.ptr<!s32i, target_address_space(4)>, !s32i) -> !cir.ptr<!s32i, target_address_space(4)> +// CIR: cir.store {{.*}} %[[ADD_PTR]], %[[APTR_ASCAST]] +// CIR: cir.scope { +// CIR: %[[APTR_VAL:.*]] = cir.load {{.*}} %[[APTR_ASCAST]] +// CIR: %[[ARR_DECAY2:.*]] = cir.cast array_to_ptrdecay %[[ARR_ASCAST]] +// CIR: %[[ADD_PTR2:.*]] = cir.ptr_stride %[[ARR_DECAY2]], {{.*}} : (!cir.ptr<!s32i, target_address_space(4)>, !u64i) -> !cir.ptr<!s32i, target_address_space(4)> +// CIR: %[[CMP2:.*]] = cir.cmp lt %[[APTR_VAL]], %[[ADD_PTR2]] : !cir.ptr<!s32i, target_address_space(4)> +// CIR: cir.if %[[CMP2]] { +// CIR: %[[APTR_VAL2:.*]] = cir.load deref {{.*}} %[[APTR_ASCAST]] +// CIR: cir.store {{.*}}, %[[APTR_VAL2]] : !s32i, !cir.ptr<!s32i, target_address_space(4)> + +// LLVM-LABEL: define {{.*}}spir_func void @_Z4testv( +// LLVM: %[[I:.*]] = alloca i32, align 4 +// LLVM-NEXT: %[[PPTR:.*]] = alloca ptr addrspace(4), align 8 +// LLVM-NEXT: %[[IS_I_PTR:.*]] = alloca i8, align 1 +// LLVM-NEXT: %[[VAR23:.*]] = alloca i32, align 4 +// LLVM-NEXT: %[[CP:.*]] = alloca ptr addrspace(4), align 8 +// LLVM-NEXT: %[[ARR:.*]] = alloca [42 x i32], align 4 +// LLVM-NEXT: %[[CPP:.*]] = alloca ptr addrspace(4), align 8 +// LLVM-NEXT: %[[APTR:.*]] = alloca ptr addrspace(4), align 8 +// LLVM-NEXT: %[[APTR_ASCAST:.*]] = addrspacecast ptr %[[APTR]] to ptr addrspace(4) +// LLVM-NEXT: %[[CPP_ASCAST:.*]] = addrspacecast ptr %[[CPP]] to ptr addrspace(4) +// LLVM-NEXT: %[[ARR_ASCAST:.*]] = addrspacecast ptr %[[ARR]] to ptr addrspace(4) +// LLVM-NEXT: %[[CP_ASCAST:.*]] = addrspacecast ptr %[[CP]] to ptr addrspace(4) +// LLVM-NEXT: %[[VAR23_ASCAST:.*]] = addrspacecast ptr %[[VAR23]] to ptr addrspace(4) +// LLVM-NEXT: %[[IS_I_PTR_ASCAST:.*]] = addrspacecast ptr %[[IS_I_PTR]] to ptr addrspace(4) +// LLVM-NEXT: %[[PPTR_ASCAST:.*]] = addrspacecast ptr %[[PPTR]] to ptr addrspace(4) +// LLVM-NEXT: %[[I_ASCAST:.*]] = addrspacecast ptr %[[I]] to ptr addrspace(4) +// LLVM-NEXT: store i32 0, ptr addrspace(4) %[[I_ASCAST]], align 4 +// LLVM-NEXT: store ptr addrspace(4) %[[I_ASCAST]], ptr addrspace(4) %[[PPTR_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP0:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[PPTR_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP1:.*]] = icmp eq ptr addrspace(4) %[[TMP0]], %[[I_ASCAST]] +// LLVM-NEXT: %[[TMP2:.*]] = zext i1 %[[TMP1]] to i8 +// LLVM-NEXT: store i8 %[[TMP2]], ptr addrspace(4) %[[IS_I_PTR_ASCAST]], align 1 +// LLVM-NEXT: store i32 23, ptr addrspace(4) %[[VAR23_ASCAST]], align 4 +// LLVM-NEXT: store ptr addrspace(4) %[[VAR23_ASCAST]], ptr addrspace(4) %[[CP_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP3:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[CP_ASCAST]], align 8 +// LLVM-NEXT: store i8 41, ptr addrspace(4) %[[TMP3]], align 1 +// LLVM-NEXT: %[[TMP4:.*]] = getelementptr i32, ptr addrspace(4) %[[ARR_ASCAST]], i32 0 +// LLVM-NEXT: store ptr addrspace(4) %[[TMP4]], ptr addrspace(4) %[[CPP_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP5:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[CPP_ASCAST]], align 8 +// LLVM-NEXT: store i8 43, ptr addrspace(4) %[[TMP5]], align 1 +// LLVM-NEXT: %[[TMP6:.*]] = getelementptr i32, ptr addrspace(4) %[[ARR_ASCAST]], i32 0 +// LLVM-NEXT: %[[TMP7:.*]] = getelementptr i32, ptr addrspace(4) %[[TMP6]], i64 10 +// LLVM-NEXT: store ptr addrspace(4) %[[TMP7]], ptr addrspace(4) %[[APTR_ASCAST]], align 8 +// LLVM-NEXT: br label %[[TMP8:.*]] +// LLVM: [[TMP8]]: +// LLVM-NEXT: %[[TMP9:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[APTR_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP10:.*]] = getelementptr i32, ptr addrspace(4) %[[ARR_ASCAST]], i32 0 +// LLVM-NEXT: %[[TMP11:.*]] = getelementptr i32, ptr addrspace(4) %[[TMP10]], i64 168 +// LLVM-NEXT: %[[TMP12:.*]] = icmp ult ptr addrspace(4) %[[TMP9]], %[[TMP11]] +// LLVM-NEXT: br i1 %[[TMP12]], label %[[TMP13:.*]], label %[[TMP14:.*]] +// LLVM: [[TMP13]]: +// LLVM-NEXT: %[[TMP15:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[APTR_ASCAST]], align 8 +// LLVM-NEXT: store i32 44, ptr addrspace(4) %[[TMP15]], align 4 +// LLVM-NEXT: br label %[[TMP14]] +// LLVM: [[TMP14]]: +// LLVM-NEXT: br label %[[TMP16:.*]] +// LLVM: [[TMP16]]: +// LLVM-NEXT: ret void + +// OGCG-LABEL: define {{.*}}spir_func void @_Z4testv( +// OGCG: %[[I:.*]] = alloca i32, align 4 +// OGCG-NEXT: %[[PPTR:.*]] = alloca ptr addrspace(4), align 8 +// OGCG-NEXT: %[[IS_I_PTR:.*]] = alloca i8, align 1 +// OGCG-NEXT: %[[VAR23:.*]] = alloca i32, align 4 +// OGCG-NEXT: %[[CP:.*]] = alloca ptr addrspace(4), align 8 +// OGCG-NEXT: %[[ARR:.*]] = alloca [42 x i32], align 4 +// OGCG-NEXT: %[[CPP:.*]] = alloca ptr addrspace(4), align 8 +// OGCG-NEXT: %[[APTR:.*]] = alloca ptr addrspace(4), align 8 +// OGCG-NEXT: %[[I_ASCAST:.*]] = addrspacecast ptr %[[I]] to ptr addrspace(4) +// OGCG-NEXT: %[[PPTR_ASCAST:.*]] = addrspacecast ptr %[[PPTR]] to ptr addrspace(4) +// OGCG-NEXT: %[[IS_I_PTR_ASCAST:.*]] = addrspacecast ptr %[[IS_I_PTR]] to ptr addrspace(4) +// OGCG-NEXT: %[[VAR23_ASCAST:.*]] = addrspacecast ptr %[[VAR23]] to ptr addrspace(4) +// OGCG-NEXT: %[[CP_ASCAST:.*]] = addrspacecast ptr %[[CP]] to ptr addrspace(4) +// OGCG-NEXT: %[[ARR_ASCAST:.*]] = addrspacecast ptr %[[ARR]] to ptr addrspace(4) +// OGCG-NEXT: %[[CPP_ASCAST:.*]] = addrspacecast ptr %[[CPP]] to ptr addrspace(4) +// OGCG-NEXT: %[[APTR_ASCAST:.*]] = addrspacecast ptr %[[APTR]] to ptr addrspace(4) +// OGCG-NEXT: store i32 0, ptr addrspace(4) %[[I_ASCAST]], align 4 +// OGCG-NEXT: store ptr addrspace(4) %[[I_ASCAST]], ptr addrspace(4) %[[PPTR_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP0:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[PPTR_ASCAST]], align 8 +// OGCG-NEXT: %[[CMP:.*]] = icmp eq ptr addrspace(4) %[[TMP0]], %[[I_ASCAST]] +// OGCG-NEXT: %[[STOREDV:.*]] = zext i1 %[[CMP]] to i8 +// OGCG-NEXT: store i8 %[[STOREDV]], ptr addrspace(4) %[[IS_I_PTR_ASCAST]], align 1 +// OGCG-NEXT: store i32 23, ptr addrspace(4) %[[VAR23_ASCAST]], align 4 +// OGCG-NEXT: store ptr addrspace(4) %[[VAR23_ASCAST]], ptr addrspace(4) %[[CP_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP1:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[CP_ASCAST]], align 8 +// OGCG-NEXT: store i8 41, ptr addrspace(4) %[[TMP1]], align 1 +// OGCG-NEXT: %[[ARRAYDECAY:.*]] = getelementptr inbounds [42 x i32], ptr addrspace(4) %[[ARR_ASCAST]], i64 0, i64 0 +// OGCG-NEXT: store ptr addrspace(4) %[[ARRAYDECAY]], ptr addrspace(4) %[[CPP_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP2:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[CPP_ASCAST]], align 8 +// OGCG-NEXT: store i8 43, ptr addrspace(4) %[[TMP2]], align 1 +// OGCG-NEXT: %[[ARRAYDECAY1:.*]] = getelementptr inbounds [42 x i32], ptr addrspace(4) %[[ARR_ASCAST]], i64 0, i64 0 +// OGCG-NEXT: %[[ADD_PTR:.*]] = getelementptr inbounds i32, ptr addrspace(4) %[[ARRAYDECAY1]], i64 10 +// OGCG-NEXT: store ptr addrspace(4) %[[ADD_PTR]], ptr addrspace(4) %[[APTR_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP3:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[APTR_ASCAST]], align 8 +// OGCG-NEXT: %[[ARRAYDECAY2:.*]] = getelementptr inbounds [42 x i32], ptr addrspace(4) %[[ARR_ASCAST]], i64 0, i64 0 +// OGCG-NEXT: %[[ADD_PTR3:.*]] = getelementptr inbounds nuw i32, ptr addrspace(4) %[[ARRAYDECAY2]], i64 168 +// OGCG-NEXT: %[[CMP4:.*]] = icmp ult ptr addrspace(4) %[[TMP3]], %[[ADD_PTR3]] +// OGCG-NEXT: br i1 %[[CMP4]], label %[[IF_THEN:.*]], label %[[IF_END:.*]] +// OGCG: [[IF_THEN]]: +// OGCG-NEXT: %[[TMP4:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[APTR_ASCAST]], align 8 +// OGCG-NEXT: store i32 44, ptr addrspace(4) %[[TMP4]], align 4 +// OGCG-NEXT: br label %[[IF_END]] +// OGCG: [[IF_END]]: +// OGCG-NEXT: ret void diff --git a/clang/test/CIR/CodeGenSYCL/local-temporaries.cpp b/clang/test/CIR/CodeGenSYCL/local-temporaries.cpp new file mode 100644 index 0000000000000..0c4649d075176 --- /dev/null +++ b/clang/test/CIR/CodeGenSYCL/local-temporaries.cpp @@ -0,0 +1,229 @@ +// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -disable-llvm-passes -std=c++17 -Wno-c99-extensions -fclangir -emit-cir %s -o %t.cir +// RUN: FileCheck --input-file=%t.cir %s -check-prefix=CIR +// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -disable-llvm-passes -std=c++17 -Wno-c99-extensions -fclangir -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --input-file=%t-cir.ll %s -check-prefix=LLVM +// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -disable-llvm-passes -std=c++17 -Wno-c99-extensions -emit-llvm %s -o %t.ll +// RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG + +// Temporaries in SYCL device code are allocated in the private address space +// and cast to the generic address space. There is no classic CodeGenSYCL +// counterpart for these cases. + +void use(int *p); +void use_const(const int *p); + +struct S { + int a; + int b; +}; +void take_ref(const S &s); + +// Temporary materialized for a reference binding. +[[clang::sycl_external]] void ref_temp() { + const int &r = 42; + use_const(&r); +} + +// Aggregate temporary passed by reference. +[[clang::sycl_external]] void agg_temp() { take_ref(S{1, 2}); } + +// Array temporary that decays to a pointer. +[[clang::sycl_external]] void array_temp() { use((int[]){3, 4}); } + +// Local captured by reference in a lambda. +[[clang::sycl_external]] void lambda_ref() { + int x = 0; + auto l = [&] { use(&x); }; + l(); +} + +// NRVO variable lives in the return slot, cast to the generic address space. +[[clang::sycl_external]] S nrvo() { + S s{}; + use(&s.a); + return s; +} + +// CIR-LABEL: cir.func {{.*}}@_Z8ref_tempv( +// CIR-NEXT: %[[REF_TMP0:.*]] = cir.alloca "ref.tmp0" align(4) init : !cir.ptr<!s32i> +// CIR-NEXT: %[[REF_TMP0_ASCAST:.*]] = cir.cast address_space %[[REF_TMP0]] : !cir.ptr<!s32i> -> !cir.ptr<!s32i, target_address_space(4)> +// CIR-NEXT: %[[R:.*]] = cir.alloca "r" align(8) init const : !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>> +// CIR-NEXT: %[[R_ASCAST:.*]] = cir.cast address_space %[[R]] : !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>> -> !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>, target_address_space(4)> +// CIR-NEXT: %[[TMP0:.*]] = cir.const #cir.int<42> : !s32i +// CIR-NEXT: cir.store align(4) %[[TMP0]], %[[REF_TMP0_ASCAST]] : !s32i, !cir.ptr<!s32i, target_address_space(4)> +// CIR-NEXT: cir.store align(8) %[[REF_TMP0_ASCAST]], %[[R_ASCAST]] : !cir.ptr<!s32i, target_address_space(4)>, !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>, target_address_space(4)> +// CIR-NEXT: %[[TMP1:.*]] = cir.load %[[R_ASCAST]] : !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>, target_address_space(4)>, !cir.ptr<!s32i, target_address_space(4)> +// CIR-NEXT: cir.call @_Z9use_constPKi(%[[TMP1]]) nothrow nounwind {convergent} : (!cir.ptr<!s32i, target_address_space(4)> {llvm.noundef}) -> () +// CIR-NEXT: cir.return + +// CIR-LABEL: cir.func {{.*}}@_Z8agg_tempv( +// CIR-NEXT: %[[REF_TMP0:.*]] = cir.alloca "ref.tmp0" align(4) : !cir.ptr<!rec_S> +// CIR-NEXT: %[[REF_TMP0_ASCAST:.*]] = cir.cast address_space %[[REF_TMP0]] : !cir.ptr<!rec_S> -> !cir.ptr<!rec_S, target_address_space(4)> +// CIR-NEXT: %[[TMP0:.*]] = cir.get_member %[[REF_TMP0_ASCAST]][0] {name = "a"} : !cir.ptr<!rec_S, target_address_space(4)> -> !cir.ptr<!s32i, target_address_space(4)> +// CIR-NEXT: %[[TMP1:.*]] = cir.const #cir.int<1> : !s32i +// CIR-NEXT: cir.store align(4) %[[TMP1]], %[[TMP0]] : !s32i, !cir.ptr<!s32i, target_address_space(4)> +// CIR-NEXT: %[[TMP2:.*]] = cir.get_member %[[REF_TMP0_ASCAST]][1] {name = "b"} : !cir.ptr<!rec_S, target_address_space(4)> -> !cir.ptr<!s32i, target_address_space(4)> +// CIR-NEXT: %[[TMP3:.*]] = cir.const #cir.int<2> : !s32i +// CIR-NEXT: cir.store align(4) %[[TMP3]], %[[TMP2]] : !s32i, !cir.ptr<!s32i, target_address_space(4)> +// CIR-NEXT: cir.call @_Z8take_refRK1S(%[[REF_TMP0_ASCAST]]) nothrow nounwind {convergent} : (!cir.ptr<!rec_S, target_address_space(4)> {llvm.align = 4 : i64, llvm.dereferenceable = 8 : i64, llvm.noundef}) -> () +// CIR-NEXT: cir.return + +// CIR-LABEL: cir.func {{.*}}@_Z10array_tempv( +// CIR-NEXT: %[[REF_TMP0:.*]] = cir.alloca "ref.tmp0" align(4) : !cir.ptr<!cir.array<!s32i x 2>> +// CIR-NEXT: %[[REF_TMP0_ASCAST:.*]] = cir.cast address_space %[[REF_TMP0]] : !cir.ptr<!cir.array<!s32i x 2>> -> !cir.ptr<!cir.array<!s32i x 2>, target_address_space(4)> +// CIR-NEXT: %[[TMP0:.*]] = cir.cast array_to_ptrdecay %[[REF_TMP0_ASCAST]] : !cir.ptr<!cir.array<!s32i x 2>, target_address_space(4)> -> !cir.ptr<!s32i, target_address_space(4)> +// CIR-NEXT: %[[TMP1:.*]] = cir.const #cir.int<3> : !s32i +// CIR-NEXT: cir.store align(4) %[[TMP1]], %[[TMP0]] : !s32i, !cir.ptr<!s32i, target_address_space(4)> +// CIR-NEXT: %[[TMP2:.*]] = cir.const #cir.int<1> : !s64i +// CIR-NEXT: %[[TMP3:.*]] = cir.ptr_stride %[[TMP0]], %[[TMP2]] : (!cir.ptr<!s32i, target_address_space(4)>, !s64i) -> !cir.ptr<!s32i, target_address_space(4)> +// CIR-NEXT: %[[TMP4:.*]] = cir.const #cir.int<4> : !s32i +// CIR-NEXT: cir.store align(4) %[[TMP4]], %[[TMP3]] : !s32i, !cir.ptr<!s32i, target_address_space(4)> +// CIR-NEXT: %[[TMP5:.*]] = cir.cast array_to_ptrdecay %[[REF_TMP0_ASCAST]] : !cir.ptr<!cir.array<!s32i x 2>, target_address_space(4)> -> !cir.ptr<!s32i, target_address_space(4)> +// CIR-NEXT: cir.call @_Z3usePi(%[[TMP5]]) nothrow nounwind {convergent} : (!cir.ptr<!s32i, target_address_space(4)> {llvm.noundef}) -> () +// CIR-NEXT: cir.return + +// CIR-LABEL: cir.func {{.*}}@_Z10lambda_refv( +// CIR-NEXT: %[[X:.*]] = cir.alloca "x" align(4) init : !cir.ptr<!s32i> +// CIR-NEXT: %[[L:.*]] = cir.alloca "l" align(8) init : !cir.ptr<!rec_anon2E0> +// CIR-NEXT: %[[L_ASCAST:.*]] = cir.cast address_space %[[L]] : !cir.ptr<!rec_anon2E0> -> !cir.ptr<!rec_anon2E0, target_address_space(4)> +// CIR-NEXT: %[[X_ASCAST:.*]] = cir.cast address_space %[[X]] : !cir.ptr<!s32i> -> !cir.ptr<!s32i, target_address_space(4)> +// CIR-NEXT: %[[TMP0:.*]] = cir.const #cir.int<0> : !s32i +// CIR-NEXT: cir.store align(4) %[[TMP0]], %[[X_ASCAST]] : !s32i, !cir.ptr<!s32i, target_address_space(4)> +// CIR-NEXT: %[[TMP1:.*]] = cir.get_member %[[L_ASCAST]][0] {name = "x"} : !cir.ptr<!rec_anon2E0, target_address_space(4)> -> !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>, target_address_space(4)> +// CIR-NEXT: cir.store align(8) %[[X_ASCAST]], %[[TMP1]] : !cir.ptr<!s32i, target_address_space(4)>, !cir.ptr<!cir.ptr<!s32i, target_address_space(4)>, target_address_space(4)> +// CIR-NEXT: cir.call @_ZZ10lambda_refvENKUlvE_clEv(%[[L_ASCAST]]) nothrow nounwind {convergent} : (!cir.ptr<!rec_anon2E0, target_address_space(4)> {llvm.align = 8 : i64, llvm.dereferenceable_or_null = 8 : i64, llvm.noundef}) -> () +// CIR-NEXT: cir.return + +// CIR-LABEL: cir.func {{.*}}@_Z4nrvov( +// CIR-NEXT: %[[__RETVAL:.*]] = cir.alloca "__retval" align(4) init : !cir.ptr<!rec_S> +// CIR-NEXT: %[[__RETVAL_ASCAST:.*]] = cir.cast address_space %[[__RETVAL]] : !cir.ptr<!rec_S> -> !cir.ptr<!rec_S, target_address_space(4)> +// CIR-NEXT: %[[TMP0:.*]] = cir.const #cir.zero : !rec_S +// CIR-NEXT: cir.store align(4) %[[TMP0]], %[[__RETVAL_ASCAST]] : !rec_S, !cir.ptr<!rec_S, target_address_space(4)> +// CIR-NEXT: %[[TMP1:.*]] = cir.get_member %[[__RETVAL_ASCAST]][0] {name = "a"} : !cir.ptr<!rec_S, target_address_space(4)> -> !cir.ptr<!s32i, target_address_space(4)> +// CIR-NEXT: cir.call @_Z3usePi(%[[TMP1]]) nothrow nounwind {convergent} : (!cir.ptr<!s32i, target_address_space(4)> {llvm.noundef}) -> () +// CIR-NEXT: %[[TMP2:.*]] = cir.load %[[__RETVAL]] : !cir.ptr<!rec_S>, !rec_S +// CIR-NEXT: cir.return %[[TMP2]] : !rec_S + +// LLVM-LABEL: define {{.*}}spir_func void @_Z8ref_tempv( +// LLVM: %[[R:.*]] = alloca i32, align 4 +// LLVM-NEXT: %[[R_ASCAST:.*]] = addrspacecast ptr %[[R]] to ptr addrspace(4) +// LLVM-NEXT: %[[REF_TMP:.*]] = alloca ptr addrspace(4), align 8 +// LLVM-NEXT: %[[REF_TMP_ASCAST:.*]] = addrspacecast ptr %[[REF_TMP]] to ptr addrspace(4) +// LLVM-NEXT: store i32 42, ptr addrspace(4) %[[R_ASCAST]], align 4 +// LLVM-NEXT: store ptr addrspace(4) %[[R_ASCAST]], ptr addrspace(4) %[[REF_TMP_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP0:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[REF_TMP_ASCAST]], align 8 +// LLVM-NEXT: call void @_Z9use_constPKi(ptr addrspace(4) noundef %[[TMP0]]) +// LLVM-NEXT: ret void + +// LLVM-LABEL: define {{.*}}spir_func void @_Z8agg_tempv( +// LLVM: %[[REF_TMP:.*]] = alloca %[[TMP0:.*]], align 4 +// LLVM-NEXT: %[[REF_TMP_ASCAST:.*]] = addrspacecast ptr %[[REF_TMP]] to ptr addrspace(4) +// LLVM-NEXT: %[[TMP1:.*]] = getelementptr inbounds nuw %[[TMP0]], ptr addrspace(4) %[[REF_TMP_ASCAST]], i32 0, i32 0 +// LLVM-NEXT: store i32 1, ptr addrspace(4) %[[TMP1]], align 4 +// LLVM-NEXT: %[[TMP2:.*]] = getelementptr inbounds nuw %[[TMP0]], ptr addrspace(4) %[[REF_TMP_ASCAST]], i32 0, i32 1 +// LLVM-NEXT: store i32 2, ptr addrspace(4) %[[TMP2]], align 4 +// LLVM-NEXT: call void @_Z8take_refRK1S(ptr addrspace(4) noundef align 4 dereferenceable(8) %[[REF_TMP_ASCAST]]) +// LLVM-NEXT: ret void + +// LLVM-LABEL: define {{.*}}spir_func void @_Z10array_tempv( +// LLVM: %[[REF_TMP:.*]] = alloca [2 x i32], align 4 +// LLVM-NEXT: %[[REF_TMP_ASCAST:.*]] = addrspacecast ptr %[[REF_TMP]] to ptr addrspace(4) +// LLVM-NEXT: %[[TMP0:.*]] = getelementptr i32, ptr addrspace(4) %[[REF_TMP_ASCAST]], i32 0 +// LLVM-NEXT: store i32 3, ptr addrspace(4) %[[TMP0]], align 4 +// LLVM-NEXT: %[[TMP1:.*]] = getelementptr i32, ptr addrspace(4) %[[TMP0]], i64 1 +// LLVM-NEXT: store i32 4, ptr addrspace(4) %[[TMP1]], align 4 +// LLVM-NEXT: %[[TMP2:.*]] = getelementptr i32, ptr addrspace(4) %[[REF_TMP_ASCAST]], i32 0 +// LLVM-NEXT: call void @_Z3usePi(ptr addrspace(4) noundef %[[TMP2]]) +// LLVM-NEXT: ret void + +// LLVM-LABEL: define {{.*}}spir_func void @_Z10lambda_refv( +// LLVM: %[[X:.*]] = alloca i32, align 4 +// LLVM-NEXT: %[[L:.*]] = alloca %[[TMP0:.*]], align 8 +// LLVM-NEXT: %[[L_ASCAST:.*]] = addrspacecast ptr %[[L]] to ptr addrspace(4) +// LLVM-NEXT: %[[X_ASCAST:.*]] = addrspacecast ptr %[[X]] to ptr addrspace(4) +// LLVM-NEXT: store i32 0, ptr addrspace(4) %[[X_ASCAST]], align 4 +// LLVM-NEXT: %[[TMP1:.*]] = getelementptr inbounds nuw %[[TMP0]], ptr addrspace(4) %[[L_ASCAST]], i32 0, i32 0 +// LLVM-NEXT: store ptr addrspace(4) %[[X_ASCAST]], ptr addrspace(4) %[[TMP1]], align 8 +// LLVM-NEXT: call void @_ZZ10lambda_refvENKUlvE_clEv(ptr addrspace(4) noundef align 8 dereferenceable_or_null(8) %[[L_ASCAST]]) +// LLVM-NEXT: ret void + +// LLVM-LABEL: define {{.*}}spir_func void @_ZZ10lambda_refvENKUlvE_clEv( +// LLVM-SAME: ptr addrspace(4) noundef align 8 dereferenceable_or_null(8) %[[THIS:.*]]) +// LLVM: %[[THIS_ADDR:.*]] = alloca ptr addrspace(4), align 8 +// LLVM-NEXT: %[[THIS_ADDR_ASCAST:.*]] = addrspacecast ptr %[[THIS_ADDR]] to ptr addrspace(4) +// LLVM-NEXT: store ptr addrspace(4) %[[THIS]], ptr addrspace(4) %[[THIS_ADDR_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP0:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[THIS_ADDR_ASCAST]], align 8 +// LLVM-NEXT: %[[TMP1:.*]] = getelementptr inbounds nuw %[[TMP2:.*]], ptr addrspace(4) %[[TMP0]], i32 0, i32 0 +// LLVM-NEXT: %[[TMP3:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[TMP1]], align 8 +// LLVM-NEXT: call void @_Z3usePi(ptr addrspace(4) noundef %[[TMP3]]) +// LLVM-NEXT: ret void + +// LLVM-LABEL: define {{.*}}spir_func %struct.S @_Z4nrvov( +// LLVM: %[[TMP0:.*]] = alloca %[[TMP1:.*]], align 4 +// LLVM-NEXT: %[[TMP2:.*]] = addrspacecast ptr %[[TMP0]] to ptr addrspace(4) +// LLVM-NEXT: store %[[TMP1]] zeroinitializer, ptr addrspace(4) %[[TMP2]], align 4 +// LLVM-NEXT: %[[TMP3:.*]] = getelementptr inbounds nuw %[[TMP1]], ptr addrspace(4) %[[TMP2]], i32 0, i32 0 +// LLVM-NEXT: call void @_Z3usePi(ptr addrspace(4) noundef %[[TMP3]]) +// LLVM-NEXT: %[[TMP4:.*]] = load %[[TMP1]], ptr %[[TMP0]], align 4 +// LLVM-NEXT: ret %[[TMP1]] %[[TMP4]] + +// OGCG-LABEL: define {{.*}}spir_func void @_Z8ref_tempv( +// OGCG: %[[R:.*]] = alloca ptr addrspace(4), align 8 +// OGCG-NEXT: %[[REF_TMP:.*]] = alloca i32, align 4 +// OGCG-NEXT: %[[R_ASCAST:.*]] = addrspacecast ptr %[[R]] to ptr addrspace(4) +// OGCG-NEXT: %[[REF_TMP_ASCAST:.*]] = addrspacecast ptr %[[REF_TMP]] to ptr addrspace(4) +// OGCG-NEXT: store i32 42, ptr addrspace(4) %[[REF_TMP_ASCAST]], align 4 +// OGCG-NEXT: store ptr addrspace(4) %[[REF_TMP_ASCAST]], ptr addrspace(4) %[[R_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP0:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[R_ASCAST]], align 8, !align !1 +// OGCG-NEXT: call spir_func void @_Z9use_constPKi(ptr addrspace(4) noundef %[[TMP0]]) +// OGCG-NEXT: ret void + +// OGCG-LABEL: define {{.*}}spir_func void @_Z8agg_tempv( +// OGCG: %[[REF_TMP:.*]] = alloca %[[STRUCT_S:.*]], align 4 +// OGCG-NEXT: %[[REF_TMP_ASCAST:.*]] = addrspacecast ptr %[[REF_TMP]] to ptr addrspace(4) +// OGCG-NEXT: %[[A:.*]] = getelementptr inbounds nuw %[[STRUCT_S]], ptr addrspace(4) %[[REF_TMP_ASCAST]], i32 0, i32 0 +// OGCG-NEXT: store i32 1, ptr addrspace(4) %[[A]], align 4 +// OGCG-NEXT: %[[B:.*]] = getelementptr inbounds nuw %[[STRUCT_S]], ptr addrspace(4) %[[REF_TMP_ASCAST]], i32 0, i32 1 +// OGCG-NEXT: store i32 2, ptr addrspace(4) %[[B]], align 4 +// OGCG-NEXT: call spir_func void @_Z8take_refRK1S(ptr addrspace(4) noundef align 4 dereferenceable(8) %[[REF_TMP_ASCAST]]) +// OGCG-NEXT: ret void + +// OGCG-LABEL: define {{.*}}spir_func void @_Z10array_tempv( +// OGCG: %[[REF_TMP:.*]] = alloca [2 x i32], align 4 +// OGCG-NEXT: %[[REF_TMP_ASCAST:.*]] = addrspacecast ptr %[[REF_TMP]] to ptr addrspace(4) +// OGCG-NEXT: store i32 3, ptr addrspace(4) %[[REF_TMP_ASCAST]], align 4 +// OGCG-NEXT: %[[ARRAYINIT_ELEMENT:.*]] = getelementptr inbounds i32, ptr addrspace(4) %[[REF_TMP_ASCAST]], i64 1 +// OGCG-NEXT: store i32 4, ptr addrspace(4) %[[ARRAYINIT_ELEMENT]], align 4 +// OGCG-NEXT: %[[ARRAYDECAY:.*]] = getelementptr inbounds [2 x i32], ptr addrspace(4) %[[REF_TMP_ASCAST]], i64 0, i64 0 +// OGCG-NEXT: call spir_func void @_Z3usePi(ptr addrspace(4) noundef %[[ARRAYDECAY]]) +// OGCG-NEXT: ret void + +// OGCG-LABEL: define {{.*}}spir_func void @_Z10lambda_refv( +// OGCG: %[[X:.*]] = alloca i32, align 4 +// OGCG-NEXT: %[[L:.*]] = alloca %[[CLASS_ANON:.*]], align 8 +// OGCG-NEXT: %[[X_ASCAST:.*]] = addrspacecast ptr %[[X]] to ptr addrspace(4) +// OGCG-NEXT: %[[L_ASCAST:.*]] = addrspacecast ptr %[[L]] to ptr addrspace(4) +// OGCG-NEXT: store i32 0, ptr addrspace(4) %[[X_ASCAST]], align 4 +// OGCG-NEXT: %[[TMP0:.*]] = getelementptr inbounds nuw %[[CLASS_ANON]], ptr addrspace(4) %[[L_ASCAST]], i32 0, i32 0 +// OGCG-NEXT: store ptr addrspace(4) %[[X_ASCAST]], ptr addrspace(4) %[[TMP0]], align 8 +// OGCG-NEXT: call spir_func void @_ZZ10lambda_refvENKUlvE_clEv(ptr addrspace(4) noundef align 8 dereferenceable_or_null(8) %[[L_ASCAST]]) +// OGCG-NEXT: ret void + +// OGCG-LABEL: define {{.*}}spir_func void @_ZZ10lambda_refvENKUlvE_clEv( +// OGCG-SAME: ptr addrspace(4) noundef align 8 dereferenceable_or_null(8) %[[THIS:.*]]) +// OGCG: %[[THIS_ADDR:.*]] = alloca ptr addrspace(4), align 8 +// OGCG-NEXT: %[[THIS_ADDR_ASCAST:.*]] = addrspacecast ptr %[[THIS_ADDR]] to ptr addrspace(4) +// OGCG-NEXT: store ptr addrspace(4) %[[THIS]], ptr addrspace(4) %[[THIS_ADDR_ASCAST]], align 8 +// OGCG-NEXT: %[[THIS1:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[THIS_ADDR_ASCAST]], align 8 +// OGCG-NEXT: %[[TMP0:.*]] = getelementptr inbounds nuw %[[CLASS_ANON:.*]], ptr addrspace(4) %[[THIS1]], i32 0, i32 0 +// OGCG-NEXT: %[[TMP1:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[TMP0]], align 8, !align !1 +// OGCG-NEXT: call spir_func void @_Z3usePi(ptr addrspace(4) noundef %[[TMP1]]) +// OGCG-NEXT: ret void + +// OGCG-LABEL: define {{.*}}spir_func void @_Z4nrvov( +// OGCG-SAME: ptr dead_on_unwind noalias writable sret(%[[STRUCT_S:.*]]) align 4 %[[AGG_RESULT:.*]]) +// OGCG: %[[AGG_RESULT_ASCAST:.*]] = addrspacecast ptr %[[AGG_RESULT]] to ptr addrspace(4) +// OGCG-NEXT: call void @llvm.memset.p4.i64(ptr addrspace(4) align 4 %[[AGG_RESULT_ASCAST]], i8 0, i64 8, i1 false) +// OGCG-NEXT: %[[A:.*]] = getelementptr inbounds nuw %[[STRUCT_S]], ptr addrspace(4) %[[AGG_RESULT_ASCAST]], i32 0, i32 0 +// OGCG-NEXT: call spir_func void @_Z3usePi(ptr addrspace(4) noundef %[[A]]) +// OGCG-NEXT: ret void diff --git a/clang/test/CIR/CodeGenSYCL/record-address-space.cpp b/clang/test/CIR/CodeGenSYCL/record-address-space.cpp index 292aab806d942..ae881857972a8 100644 --- a/clang/test/CIR/CodeGenSYCL/record-address-space.cpp +++ b/clang/test/CIR/CodeGenSYCL/record-address-space.cpp @@ -27,7 +27,7 @@ struct S { } // POST-LABEL: cir.func {{.*}}@_Z12store_memberPU3AS1i -// POST: cir.get_member %{{.*}}[0] {name = "p"} : !cir.ptr<!rec_S> -> !cir.ptr<!cir.ptr<!s32i, target_address_space(1)>> +// POST: cir.get_member %{{.*}}[0] {name = "p"} : !cir.ptr<!rec_S, target_address_space(4)> -> !cir.ptr<!cir.ptr<!s32i, target_address_space(1)>, target_address_space(4)> // LLVM-LABEL: define {{.*}}@_Z12store_memberPU3AS1i // LLVM: %[[P:.*]] = getelementptr inbounds nuw %struct.S, ptr {{.*}}, i32 0, i32 0 >From 97fa3d10ee71534056ed5c18e40542cffb8069f2 Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris <[email protected]> Date: Mon, 5 Oct 2026 17:40:54 -0700 Subject: [PATCH 2/4] [CIR][SYCL] Share identical LLVM checks between CIR and classic CodeGen Co-Authored-By: Claude Opus 5.5 (1M context) <[email protected]> --- .../CodeGenSYCL/address-space-conversions.cpp | 350 +++++++----------- 1 file changed, 140 insertions(+), 210 deletions(-) diff --git a/clang/test/CIR/CodeGenSYCL/address-space-conversions.cpp b/clang/test/CIR/CodeGenSYCL/address-space-conversions.cpp index 4ea1e603cd87a..5a2213a8809d2 100644 --- a/clang/test/CIR/CodeGenSYCL/address-space-conversions.cpp +++ b/clang/test/CIR/CodeGenSYCL/address-space-conversions.cpp @@ -1,9 +1,9 @@ // RUN: %clang_cc1 -triple spir64 -fsycl-is-device -disable-llvm-passes -Wno-deprecated-attributes -fclangir -emit-cir %s -o %t.cir // RUN: FileCheck --input-file=%t.cir %s -check-prefix=CIR // RUN: %clang_cc1 -triple spir64 -fsycl-is-device -disable-llvm-passes -Wno-deprecated-attributes -fclangir -emit-llvm %s -o %t-cir.ll -// RUN: FileCheck --input-file=%t-cir.ll %s -check-prefix=LLVM +// RUN: FileCheck --input-file=%t-cir.ll %s --check-prefixes=LLVM,LLVM-CIR // RUN: %clang_cc1 -triple spir64 -fsycl-is-device -disable-llvm-passes -Wno-deprecated-attributes -emit-llvm %s -o %t.ll -// RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG +// RUN: FileCheck --input-file=%t.ll %s --check-prefixes=LLVM,OGCG // Port of clang/test/CodeGenSYCL/address-space-conversions.cpp. Local // variables and parameters are allocated in the private address space and @@ -109,154 +109,84 @@ void tmpl(T t) {} // CIR-NEXT: cir.store %arg0, %[[DATA_ASCAST]] : !cir.ptr<!s32i, target_address_space(3)>, !cir.ptr<!cir.ptr<!s32i, target_address_space(3)>, target_address_space(4)> // CIR-NEXT: cir.return -// LLVM-LABEL: define {{.*}}spir_func void @_Z6usagesv( -// LLVM: %[[NOAS:.*]] = alloca ptr addrspace(4), align 8 -// LLVM-NEXT: %[[GLOB:.*]] = alloca ptr addrspace(1), align 8 -// LLVM-NEXT: %[[LOC:.*]] = alloca ptr addrspace(3), align 8 -// LLVM-NEXT: %[[PRIV:.*]] = alloca ptr, align 8 -// LLVM-NEXT: %[[GLOBDEVICE:.*]] = alloca ptr addrspace(5), align 8 -// LLVM-NEXT: %[[GLOBHOST:.*]] = alloca ptr addrspace(6), align 8 -// LLVM-NEXT: %[[GLOBHOST_ASCAST:.*]] = addrspacecast ptr %[[GLOBHOST]] to ptr addrspace(4) -// LLVM-NEXT: %[[GLOBDEVICE_ASCAST:.*]] = addrspacecast ptr %[[GLOBDEVICE]] to ptr addrspace(4) -// LLVM-NEXT: %[[PRIV_ASCAST:.*]] = addrspacecast ptr %[[PRIV]] to ptr addrspace(4) -// LLVM-NEXT: %[[LOC_ASCAST:.*]] = addrspacecast ptr %[[LOC]] to ptr addrspace(4) -// LLVM-NEXT: %[[GLOB_ASCAST:.*]] = addrspacecast ptr %[[GLOB]] to ptr addrspace(4) -// LLVM-NEXT: %[[NOAS_ASCAST:.*]] = addrspacecast ptr %[[NOAS]] to ptr addrspace(4) -// LLVM-NEXT: store ptr addrspace(3) null, ptr addrspace(4) %[[LOC_ASCAST]], align 8 -// LLVM-NEXT: store ptr addrspace(1) null, ptr addrspace(4) %[[GLOB_ASCAST]], align 8 -// LLVM-NEXT: %[[TMP0:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 -// LLVM-NEXT: %[[TMP1:.*]] = addrspacecast ptr addrspace(1) %[[TMP0]] to ptr addrspace(4) -// LLVM-NEXT: store ptr addrspace(4) %[[TMP1]], ptr addrspace(4) %[[NOAS_ASCAST]], align 8 -// LLVM-NEXT: %[[TMP2:.*]] = load ptr addrspace(3), ptr addrspace(4) %[[LOC_ASCAST]], align 8 -// LLVM-NEXT: %[[TMP3:.*]] = addrspacecast ptr addrspace(3) %[[TMP2]] to ptr addrspace(4) -// LLVM-NEXT: store ptr addrspace(4) %[[TMP3]], ptr addrspace(4) %[[NOAS_ASCAST]], align 8 -// LLVM-NEXT: %[[TMP4:.*]] = load ptr, ptr addrspace(4) %[[PRIV_ASCAST]], align 8 -// LLVM-NEXT: %[[TMP5:.*]] = addrspacecast ptr %[[TMP4]] to ptr addrspace(4) -// LLVM-NEXT: store ptr addrspace(4) %[[TMP5]], ptr addrspace(4) %[[NOAS_ASCAST]], align 8 -// LLVM-NEXT: %[[TMP6:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 -// LLVM-NEXT: %[[TMP7:.*]] = addrspacecast ptr addrspace(4) %[[TMP6]] to ptr addrspace(1) -// LLVM-NEXT: store ptr addrspace(1) %[[TMP7]], ptr addrspace(4) %[[GLOB_ASCAST]], align 8 -// LLVM-NEXT: %[[TMP8:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 -// LLVM-NEXT: %[[TMP9:.*]] = addrspacecast ptr addrspace(4) %[[TMP8]] to ptr addrspace(3) -// LLVM-NEXT: store ptr addrspace(3) %[[TMP9]], ptr addrspace(4) %[[LOC_ASCAST]], align 8 -// LLVM-NEXT: %[[TMP10:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 -// LLVM-NEXT: %[[TMP11:.*]] = addrspacecast ptr addrspace(4) %[[TMP10]] to ptr -// LLVM-NEXT: store ptr %[[TMP11]], ptr addrspace(4) %[[PRIV_ASCAST]], align 8 -// LLVM-NEXT: %[[TMP12:.*]] = load ptr addrspace(5), ptr addrspace(4) %[[GLOBDEVICE_ASCAST]], align 8 -// LLVM-NEXT: %[[TMP13:.*]] = addrspacecast ptr addrspace(5) %[[TMP12]] to ptr addrspace(1) -// LLVM-NEXT: store ptr addrspace(1) %[[TMP13]], ptr addrspace(4) %[[GLOB_ASCAST]], align 8 -// LLVM-NEXT: %[[TMP14:.*]] = load ptr addrspace(6), ptr addrspace(4) %[[GLOBHOST_ASCAST]], align 8 -// LLVM-NEXT: %[[TMP15:.*]] = addrspacecast ptr addrspace(6) %[[TMP14]] to ptr addrspace(1) -// LLVM-NEXT: store ptr addrspace(1) %[[TMP15]], ptr addrspace(4) %[[GLOB_ASCAST]], align 8 -// LLVM-NEXT: %[[TMP16:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 -// LLVM-NEXT: %[[TMP17:.*]] = addrspacecast ptr addrspace(1) %[[TMP16]] to ptr addrspace(4) -// LLVM-NEXT: call void @_Z3barRi(ptr addrspace(4) noundef align 4 dereferenceable(4) %[[TMP17]]) -// LLVM-NEXT: %[[TMP18:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 -// LLVM-NEXT: %[[TMP19:.*]] = addrspacecast ptr addrspace(1) %[[TMP18]] to ptr addrspace(4) -// LLVM-NEXT: call void @_Z4bar2Ri(ptr addrspace(4) noundef align 4 dereferenceable(4) %[[TMP19]]) -// LLVM-NEXT: %[[TMP20:.*]] = load ptr addrspace(3), ptr addrspace(4) %[[LOC_ASCAST]], align 8 -// LLVM-NEXT: call void @_Z3barRU3AS3i(ptr addrspace(3) noundef align 4 dereferenceable(4) %[[TMP20]]) -// LLVM-NEXT: %[[TMP21:.*]] = load ptr addrspace(3), ptr addrspace(4) %[[LOC_ASCAST]], align 8 -// LLVM-NEXT: %[[TMP22:.*]] = addrspacecast ptr addrspace(3) %[[TMP21]] to ptr addrspace(4) -// LLVM-NEXT: call void @_Z4bar2Ri(ptr addrspace(4) noundef align 4 dereferenceable(4) %[[TMP22]]) -// LLVM-NEXT: %[[TMP23:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 -// LLVM-NEXT: call void @_Z3barRi(ptr addrspace(4) noundef align 4 dereferenceable(4) %[[TMP23]]) -// LLVM-NEXT: %[[TMP24:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 -// LLVM-NEXT: call void @_Z4bar2Ri(ptr addrspace(4) noundef align 4 dereferenceable(4) %[[TMP24]]) -// LLVM-NEXT: %[[TMP25:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 -// LLVM-NEXT: %[[TMP26:.*]] = addrspacecast ptr addrspace(1) %[[TMP25]] to ptr addrspace(4) -// LLVM-NEXT: call void @_Z3fooPi(ptr addrspace(4) noundef %[[TMP26]]) -// LLVM-NEXT: %[[TMP27:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 -// LLVM-NEXT: %[[TMP28:.*]] = addrspacecast ptr addrspace(1) %[[TMP27]] to ptr addrspace(4) -// LLVM-NEXT: call void @_Z4foo2Pi(ptr addrspace(4) noundef %[[TMP28]]) -// LLVM-NEXT: %[[TMP29:.*]] = load ptr addrspace(3), ptr addrspace(4) %[[LOC_ASCAST]], align 8 -// LLVM-NEXT: call void @_Z3fooPU3AS3i(ptr addrspace(3) noundef %[[TMP29]]) -// LLVM-NEXT: %[[TMP30:.*]] = load ptr addrspace(3), ptr addrspace(4) %[[LOC_ASCAST]], align 8 -// LLVM-NEXT: %[[TMP31:.*]] = addrspacecast ptr addrspace(3) %[[TMP30]] to ptr addrspace(4) -// LLVM-NEXT: call void @_Z4foo2Pi(ptr addrspace(4) noundef %[[TMP31]]) -// LLVM-NEXT: %[[TMP32:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 -// LLVM-NEXT: call void @_Z3fooPi(ptr addrspace(4) noundef %[[TMP32]]) -// LLVM-NEXT: %[[TMP33:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 -// LLVM-NEXT: call void @_Z4foo2Pi(ptr addrspace(4) noundef %[[TMP33]]) -// LLVM-NEXT: %[[TMP34:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 -// LLVM-NEXT: call void @_Z4tmplIPU3AS1iEvT_(ptr addrspace(1) noundef %[[TMP34]]) -// LLVM-NEXT: %[[TMP35:.*]] = load ptr addrspace(3), ptr addrspace(4) %[[LOC_ASCAST]], align 8 -// LLVM-NEXT: call void @_Z4tmplIPU3AS3iEvT_(ptr addrspace(3) noundef %[[TMP35]]) -// LLVM-NEXT: %[[TMP36:.*]] = load ptr, ptr addrspace(4) %[[PRIV_ASCAST]], align 8 -// LLVM-NEXT: call void @_Z4tmplIPU3AS0iEvT_(ptr noundef %[[TMP36]]) -// LLVM-NEXT: %[[TMP37:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 -// LLVM-NEXT: call void @_Z4tmplIPiEvT_(ptr addrspace(4) noundef %[[TMP37]]) -// LLVM-NEXT: ret void - -// LLVM-LABEL: define {{.*}}spir_func void @_Z3barRi( -// LLVM-SAME: ptr addrspace(4) noundef align 4 dereferenceable(4) %[[DATA:.*]]) -// LLVM: %[[DATA_ADDR:.*]] = alloca ptr addrspace(4), align 8 -// LLVM-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) -// LLVM-NEXT: store ptr addrspace(4) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 -// LLVM-NEXT: ret void - -// LLVM-LABEL: define {{.*}}spir_func void @_Z4bar2Ri( -// LLVM-SAME: ptr addrspace(4) noundef align 4 dereferenceable(4) %[[DATA:.*]]) -// LLVM: %[[DATA_ADDR:.*]] = alloca ptr addrspace(4), align 8 -// LLVM-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) -// LLVM-NEXT: store ptr addrspace(4) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 -// LLVM-NEXT: ret void - -// LLVM-LABEL: define {{.*}}spir_func void @_Z3barRU3AS3i( -// LLVM-SAME: ptr addrspace(3) noundef align 4 dereferenceable(4) %[[DATA:.*]]) -// LLVM: %[[DATA_ADDR:.*]] = alloca ptr addrspace(3), align 8 -// LLVM-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) -// LLVM-NEXT: store ptr addrspace(3) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 -// LLVM-NEXT: ret void - -// LLVM-LABEL: define {{.*}}spir_func void @_Z3fooPi( -// LLVM-SAME: ptr addrspace(4) noundef %[[DATA:.*]]) -// LLVM: %[[DATA_ADDR:.*]] = alloca ptr addrspace(4), align 8 -// LLVM-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) -// LLVM-NEXT: store ptr addrspace(4) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 -// LLVM-NEXT: ret void - -// LLVM-LABEL: define {{.*}}spir_func void @_Z4foo2Pi( -// LLVM-SAME: ptr addrspace(4) noundef %[[DATA:.*]]) -// LLVM: %[[DATA_ADDR:.*]] = alloca ptr addrspace(4), align 8 -// LLVM-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) -// LLVM-NEXT: store ptr addrspace(4) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 -// LLVM-NEXT: ret void - -// LLVM-LABEL: define {{.*}}spir_func void @_Z3fooPU3AS3i( -// LLVM-SAME: ptr addrspace(3) noundef %[[DATA:.*]]) -// LLVM: %[[DATA_ADDR:.*]] = alloca ptr addrspace(3), align 8 -// LLVM-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) -// LLVM-NEXT: store ptr addrspace(3) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 -// LLVM-NEXT: ret void - -// LLVM-LABEL: define {{.*}}spir_func void @_Z4tmplIPU3AS1iEvT_( -// LLVM-SAME: ptr addrspace(1) noundef %[[T:.*]]) -// LLVM: %[[T_ADDR:.*]] = alloca ptr addrspace(1), align 8 -// LLVM-NEXT: %[[T_ADDR_ASCAST:.*]] = addrspacecast ptr %[[T_ADDR]] to ptr addrspace(4) -// LLVM-NEXT: store ptr addrspace(1) %[[T]], ptr addrspace(4) %[[T_ADDR_ASCAST]], align 8 -// LLVM-NEXT: ret void - -// LLVM-LABEL: define {{.*}}spir_func void @_Z4tmplIPU3AS3iEvT_( -// LLVM-SAME: ptr addrspace(3) noundef %[[T:.*]]) -// LLVM: %[[T_ADDR:.*]] = alloca ptr addrspace(3), align 8 -// LLVM-NEXT: %[[T_ADDR_ASCAST:.*]] = addrspacecast ptr %[[T_ADDR]] to ptr addrspace(4) -// LLVM-NEXT: store ptr addrspace(3) %[[T]], ptr addrspace(4) %[[T_ADDR_ASCAST]], align 8 -// LLVM-NEXT: ret void - -// LLVM-LABEL: define {{.*}}spir_func void @_Z4tmplIPU3AS0iEvT_( -// LLVM-SAME: ptr noundef %[[T:.*]]) -// LLVM: %[[T_ADDR:.*]] = alloca ptr, align 8 -// LLVM-NEXT: %[[T_ADDR_ASCAST:.*]] = addrspacecast ptr %[[T_ADDR]] to ptr addrspace(4) -// LLVM-NEXT: store ptr %[[T]], ptr addrspace(4) %[[T_ADDR_ASCAST]], align 8 -// LLVM-NEXT: ret void - -// LLVM-LABEL: define {{.*}}spir_func void @_Z4tmplIPiEvT_( -// LLVM-SAME: ptr addrspace(4) noundef %[[T:.*]]) -// LLVM: %[[T_ADDR:.*]] = alloca ptr addrspace(4), align 8 -// LLVM-NEXT: %[[T_ADDR_ASCAST:.*]] = addrspacecast ptr %[[T_ADDR]] to ptr addrspace(4) -// LLVM-NEXT: store ptr addrspace(4) %[[T]], ptr addrspace(4) %[[T_ADDR_ASCAST]], align 8 -// LLVM-NEXT: ret void +// LLVM-CIR-LABEL: define {{.*}}spir_func void @_Z6usagesv( +// LLVM-CIR: %[[NOAS:.*]] = alloca ptr addrspace(4), align 8 +// LLVM-CIR-NEXT: %[[GLOB:.*]] = alloca ptr addrspace(1), align 8 +// LLVM-CIR-NEXT: %[[LOC:.*]] = alloca ptr addrspace(3), align 8 +// LLVM-CIR-NEXT: %[[PRIV:.*]] = alloca ptr, align 8 +// LLVM-CIR-NEXT: %[[GLOBDEVICE:.*]] = alloca ptr addrspace(5), align 8 +// LLVM-CIR-NEXT: %[[GLOBHOST:.*]] = alloca ptr addrspace(6), align 8 +// LLVM-CIR-NEXT: %[[GLOBHOST_ASCAST:.*]] = addrspacecast ptr %[[GLOBHOST]] to ptr addrspace(4) +// LLVM-CIR-NEXT: %[[GLOBDEVICE_ASCAST:.*]] = addrspacecast ptr %[[GLOBDEVICE]] to ptr addrspace(4) +// LLVM-CIR-NEXT: %[[PRIV_ASCAST:.*]] = addrspacecast ptr %[[PRIV]] to ptr addrspace(4) +// LLVM-CIR-NEXT: %[[LOC_ASCAST:.*]] = addrspacecast ptr %[[LOC]] to ptr addrspace(4) +// LLVM-CIR-NEXT: %[[GLOB_ASCAST:.*]] = addrspacecast ptr %[[GLOB]] to ptr addrspace(4) +// LLVM-CIR-NEXT: %[[NOAS_ASCAST:.*]] = addrspacecast ptr %[[NOAS]] to ptr addrspace(4) +// LLVM-CIR-NEXT: store ptr addrspace(3) null, ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// LLVM-CIR-NEXT: store ptr addrspace(1) null, ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// LLVM-CIR-NEXT: %[[TMP0:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// LLVM-CIR-NEXT: %[[TMP1:.*]] = addrspacecast ptr addrspace(1) %[[TMP0]] to ptr addrspace(4) +// LLVM-CIR-NEXT: store ptr addrspace(4) %[[TMP1]], ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// LLVM-CIR-NEXT: %[[TMP2:.*]] = load ptr addrspace(3), ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// LLVM-CIR-NEXT: %[[TMP3:.*]] = addrspacecast ptr addrspace(3) %[[TMP2]] to ptr addrspace(4) +// LLVM-CIR-NEXT: store ptr addrspace(4) %[[TMP3]], ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// LLVM-CIR-NEXT: %[[TMP4:.*]] = load ptr, ptr addrspace(4) %[[PRIV_ASCAST]], align 8 +// LLVM-CIR-NEXT: %[[TMP5:.*]] = addrspacecast ptr %[[TMP4]] to ptr addrspace(4) +// LLVM-CIR-NEXT: store ptr addrspace(4) %[[TMP5]], ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// LLVM-CIR-NEXT: %[[TMP6:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// LLVM-CIR-NEXT: %[[TMP7:.*]] = addrspacecast ptr addrspace(4) %[[TMP6]] to ptr addrspace(1) +// LLVM-CIR-NEXT: store ptr addrspace(1) %[[TMP7]], ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// LLVM-CIR-NEXT: %[[TMP8:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// LLVM-CIR-NEXT: %[[TMP9:.*]] = addrspacecast ptr addrspace(4) %[[TMP8]] to ptr addrspace(3) +// LLVM-CIR-NEXT: store ptr addrspace(3) %[[TMP9]], ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// LLVM-CIR-NEXT: %[[TMP10:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// LLVM-CIR-NEXT: %[[TMP11:.*]] = addrspacecast ptr addrspace(4) %[[TMP10]] to ptr +// LLVM-CIR-NEXT: store ptr %[[TMP11]], ptr addrspace(4) %[[PRIV_ASCAST]], align 8 +// LLVM-CIR-NEXT: %[[TMP12:.*]] = load ptr addrspace(5), ptr addrspace(4) %[[GLOBDEVICE_ASCAST]], align 8 +// LLVM-CIR-NEXT: %[[TMP13:.*]] = addrspacecast ptr addrspace(5) %[[TMP12]] to ptr addrspace(1) +// LLVM-CIR-NEXT: store ptr addrspace(1) %[[TMP13]], ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// LLVM-CIR-NEXT: %[[TMP14:.*]] = load ptr addrspace(6), ptr addrspace(4) %[[GLOBHOST_ASCAST]], align 8 +// LLVM-CIR-NEXT: %[[TMP15:.*]] = addrspacecast ptr addrspace(6) %[[TMP14]] to ptr addrspace(1) +// LLVM-CIR-NEXT: store ptr addrspace(1) %[[TMP15]], ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// LLVM-CIR-NEXT: %[[TMP16:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// LLVM-CIR-NEXT: %[[TMP17:.*]] = addrspacecast ptr addrspace(1) %[[TMP16]] to ptr addrspace(4) +// LLVM-CIR-NEXT: call void @_Z3barRi(ptr addrspace(4) noundef align 4 dereferenceable(4) %[[TMP17]]) +// LLVM-CIR-NEXT: %[[TMP18:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// LLVM-CIR-NEXT: %[[TMP19:.*]] = addrspacecast ptr addrspace(1) %[[TMP18]] to ptr addrspace(4) +// LLVM-CIR-NEXT: call void @_Z4bar2Ri(ptr addrspace(4) noundef align 4 dereferenceable(4) %[[TMP19]]) +// LLVM-CIR-NEXT: %[[TMP20:.*]] = load ptr addrspace(3), ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// LLVM-CIR-NEXT: call void @_Z3barRU3AS3i(ptr addrspace(3) noundef align 4 dereferenceable(4) %[[TMP20]]) +// LLVM-CIR-NEXT: %[[TMP21:.*]] = load ptr addrspace(3), ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// LLVM-CIR-NEXT: %[[TMP22:.*]] = addrspacecast ptr addrspace(3) %[[TMP21]] to ptr addrspace(4) +// LLVM-CIR-NEXT: call void @_Z4bar2Ri(ptr addrspace(4) noundef align 4 dereferenceable(4) %[[TMP22]]) +// LLVM-CIR-NEXT: %[[TMP23:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// LLVM-CIR-NEXT: call void @_Z3barRi(ptr addrspace(4) noundef align 4 dereferenceable(4) %[[TMP23]]) +// LLVM-CIR-NEXT: %[[TMP24:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// LLVM-CIR-NEXT: call void @_Z4bar2Ri(ptr addrspace(4) noundef align 4 dereferenceable(4) %[[TMP24]]) +// LLVM-CIR-NEXT: %[[TMP25:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// LLVM-CIR-NEXT: %[[TMP26:.*]] = addrspacecast ptr addrspace(1) %[[TMP25]] to ptr addrspace(4) +// LLVM-CIR-NEXT: call void @_Z3fooPi(ptr addrspace(4) noundef %[[TMP26]]) +// LLVM-CIR-NEXT: %[[TMP27:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// LLVM-CIR-NEXT: %[[TMP28:.*]] = addrspacecast ptr addrspace(1) %[[TMP27]] to ptr addrspace(4) +// LLVM-CIR-NEXT: call void @_Z4foo2Pi(ptr addrspace(4) noundef %[[TMP28]]) +// LLVM-CIR-NEXT: %[[TMP29:.*]] = load ptr addrspace(3), ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// LLVM-CIR-NEXT: call void @_Z3fooPU3AS3i(ptr addrspace(3) noundef %[[TMP29]]) +// LLVM-CIR-NEXT: %[[TMP30:.*]] = load ptr addrspace(3), ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// LLVM-CIR-NEXT: %[[TMP31:.*]] = addrspacecast ptr addrspace(3) %[[TMP30]] to ptr addrspace(4) +// LLVM-CIR-NEXT: call void @_Z4foo2Pi(ptr addrspace(4) noundef %[[TMP31]]) +// LLVM-CIR-NEXT: %[[TMP32:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// LLVM-CIR-NEXT: call void @_Z3fooPi(ptr addrspace(4) noundef %[[TMP32]]) +// LLVM-CIR-NEXT: %[[TMP33:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// LLVM-CIR-NEXT: call void @_Z4foo2Pi(ptr addrspace(4) noundef %[[TMP33]]) +// LLVM-CIR-NEXT: %[[TMP34:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// LLVM-CIR-NEXT: call void @_Z4tmplIPU3AS1iEvT_(ptr addrspace(1) noundef %[[TMP34]]) +// LLVM-CIR-NEXT: %[[TMP35:.*]] = load ptr addrspace(3), ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// LLVM-CIR-NEXT: call void @_Z4tmplIPU3AS3iEvT_(ptr addrspace(3) noundef %[[TMP35]]) +// LLVM-CIR-NEXT: %[[TMP36:.*]] = load ptr, ptr addrspace(4) %[[PRIV_ASCAST]], align 8 +// LLVM-CIR-NEXT: call void @_Z4tmplIPU3AS0iEvT_(ptr noundef %[[TMP36]]) +// LLVM-CIR-NEXT: %[[TMP37:.*]] = load ptr addrspace(4), ptr addrspace(4) %[[NOAS_ASCAST]], align 8 +// LLVM-CIR-NEXT: call void @_Z4tmplIPiEvT_(ptr addrspace(4) noundef %[[TMP37]]) +// LLVM-CIR-NEXT: ret void // OGCG-LABEL: define {{.*}}spir_func void @_Z6usagesv( // OGCG: %[[NOAS:.*]] = alloca ptr addrspace(4), align 8 @@ -337,72 +267,72 @@ void tmpl(T t) {} // OGCG-NEXT: call spir_func void @_Z4tmplIPiEvT_(ptr addrspace(4) noundef %[[TMP37]]) // OGCG-NEXT: ret void -// OGCG-LABEL: define {{.*}}spir_func void @_Z3barRi( -// OGCG-SAME: ptr addrspace(4) noundef align 4 dereferenceable(4) %[[DATA:.*]]) -// OGCG: %[[DATA_ADDR:.*]] = alloca ptr addrspace(4), align 8 -// OGCG-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) -// OGCG-NEXT: store ptr addrspace(4) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 -// OGCG-NEXT: ret void +// LLVM-LABEL: define {{.*}}spir_func void @_Z3barRi( +// LLVM-SAME: ptr addrspace(4) noundef align 4 dereferenceable(4) %[[DATA:.*]]) +// LLVM: %[[DATA_ADDR:.*]] = alloca ptr addrspace(4), align 8 +// LLVM-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) +// LLVM-NEXT: store ptr addrspace(4) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 +// LLVM-NEXT: ret void -// OGCG-LABEL: define {{.*}}spir_func void @_Z4bar2Ri( -// OGCG-SAME: ptr addrspace(4) noundef align 4 dereferenceable(4) %[[DATA:.*]]) -// OGCG: %[[DATA_ADDR:.*]] = alloca ptr addrspace(4), align 8 -// OGCG-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) -// OGCG-NEXT: store ptr addrspace(4) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 -// OGCG-NEXT: ret void +// LLVM-LABEL: define {{.*}}spir_func void @_Z4bar2Ri( +// LLVM-SAME: ptr addrspace(4) noundef align 4 dereferenceable(4) %[[DATA:.*]]) +// LLVM: %[[DATA_ADDR:.*]] = alloca ptr addrspace(4), align 8 +// LLVM-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) +// LLVM-NEXT: store ptr addrspace(4) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 +// LLVM-NEXT: ret void -// OGCG-LABEL: define {{.*}}spir_func void @_Z3barRU3AS3i( -// OGCG-SAME: ptr addrspace(3) noundef align 4 dereferenceable(4) %[[DATA:.*]]) -// OGCG: %[[DATA_ADDR:.*]] = alloca ptr addrspace(3), align 8 -// OGCG-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) -// OGCG-NEXT: store ptr addrspace(3) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 -// OGCG-NEXT: ret void +// LLVM-LABEL: define {{.*}}spir_func void @_Z3barRU3AS3i( +// LLVM-SAME: ptr addrspace(3) noundef align 4 dereferenceable(4) %[[DATA:.*]]) +// LLVM: %[[DATA_ADDR:.*]] = alloca ptr addrspace(3), align 8 +// LLVM-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) +// LLVM-NEXT: store ptr addrspace(3) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 +// LLVM-NEXT: ret void -// OGCG-LABEL: define {{.*}}spir_func void @_Z3fooPi( -// OGCG-SAME: ptr addrspace(4) noundef %[[DATA:.*]]) -// OGCG: %[[DATA_ADDR:.*]] = alloca ptr addrspace(4), align 8 -// OGCG-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) -// OGCG-NEXT: store ptr addrspace(4) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 -// OGCG-NEXT: ret void +// LLVM-LABEL: define {{.*}}spir_func void @_Z3fooPi( +// LLVM-SAME: ptr addrspace(4) noundef %[[DATA:.*]]) +// LLVM: %[[DATA_ADDR:.*]] = alloca ptr addrspace(4), align 8 +// LLVM-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) +// LLVM-NEXT: store ptr addrspace(4) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 +// LLVM-NEXT: ret void -// OGCG-LABEL: define {{.*}}spir_func void @_Z4foo2Pi( -// OGCG-SAME: ptr addrspace(4) noundef %[[DATA:.*]]) -// OGCG: %[[DATA_ADDR:.*]] = alloca ptr addrspace(4), align 8 -// OGCG-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) -// OGCG-NEXT: store ptr addrspace(4) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 -// OGCG-NEXT: ret void +// LLVM-LABEL: define {{.*}}spir_func void @_Z4foo2Pi( +// LLVM-SAME: ptr addrspace(4) noundef %[[DATA:.*]]) +// LLVM: %[[DATA_ADDR:.*]] = alloca ptr addrspace(4), align 8 +// LLVM-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) +// LLVM-NEXT: store ptr addrspace(4) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 +// LLVM-NEXT: ret void -// OGCG-LABEL: define {{.*}}spir_func void @_Z3fooPU3AS3i( -// OGCG-SAME: ptr addrspace(3) noundef %[[DATA:.*]]) -// OGCG: %[[DATA_ADDR:.*]] = alloca ptr addrspace(3), align 8 -// OGCG-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) -// OGCG-NEXT: store ptr addrspace(3) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 -// OGCG-NEXT: ret void +// LLVM-LABEL: define {{.*}}spir_func void @_Z3fooPU3AS3i( +// LLVM-SAME: ptr addrspace(3) noundef %[[DATA:.*]]) +// LLVM: %[[DATA_ADDR:.*]] = alloca ptr addrspace(3), align 8 +// LLVM-NEXT: %[[DATA_ADDR_ASCAST:.*]] = addrspacecast ptr %[[DATA_ADDR]] to ptr addrspace(4) +// LLVM-NEXT: store ptr addrspace(3) %[[DATA]], ptr addrspace(4) %[[DATA_ADDR_ASCAST]], align 8 +// LLVM-NEXT: ret void -// OGCG-LABEL: define {{.*}}spir_func void @_Z4tmplIPU3AS1iEvT_( -// OGCG-SAME: ptr addrspace(1) noundef %[[T:.*]]) -// OGCG: %[[T_ADDR:.*]] = alloca ptr addrspace(1), align 8 -// OGCG-NEXT: %[[T_ADDR_ASCAST:.*]] = addrspacecast ptr %[[T_ADDR]] to ptr addrspace(4) -// OGCG-NEXT: store ptr addrspace(1) %[[T]], ptr addrspace(4) %[[T_ADDR_ASCAST]], align 8 -// OGCG-NEXT: ret void +// LLVM-LABEL: define {{.*}}spir_func void @_Z4tmplIPU3AS1iEvT_( +// LLVM-SAME: ptr addrspace(1) noundef %[[T:.*]]) +// LLVM: %[[T_ADDR:.*]] = alloca ptr addrspace(1), align 8 +// LLVM-NEXT: %[[T_ADDR_ASCAST:.*]] = addrspacecast ptr %[[T_ADDR]] to ptr addrspace(4) +// LLVM-NEXT: store ptr addrspace(1) %[[T]], ptr addrspace(4) %[[T_ADDR_ASCAST]], align 8 +// LLVM-NEXT: ret void -// OGCG-LABEL: define {{.*}}spir_func void @_Z4tmplIPU3AS3iEvT_( -// OGCG-SAME: ptr addrspace(3) noundef %[[T:.*]]) -// OGCG: %[[T_ADDR:.*]] = alloca ptr addrspace(3), align 8 -// OGCG-NEXT: %[[T_ADDR_ASCAST:.*]] = addrspacecast ptr %[[T_ADDR]] to ptr addrspace(4) -// OGCG-NEXT: store ptr addrspace(3) %[[T]], ptr addrspace(4) %[[T_ADDR_ASCAST]], align 8 -// OGCG-NEXT: ret void +// LLVM-LABEL: define {{.*}}spir_func void @_Z4tmplIPU3AS3iEvT_( +// LLVM-SAME: ptr addrspace(3) noundef %[[T:.*]]) +// LLVM: %[[T_ADDR:.*]] = alloca ptr addrspace(3), align 8 +// LLVM-NEXT: %[[T_ADDR_ASCAST:.*]] = addrspacecast ptr %[[T_ADDR]] to ptr addrspace(4) +// LLVM-NEXT: store ptr addrspace(3) %[[T]], ptr addrspace(4) %[[T_ADDR_ASCAST]], align 8 +// LLVM-NEXT: ret void -// OGCG-LABEL: define {{.*}}spir_func void @_Z4tmplIPU3AS0iEvT_( -// OGCG-SAME: ptr noundef %[[T:.*]]) -// OGCG: %[[T_ADDR:.*]] = alloca ptr, align 8 -// OGCG-NEXT: %[[T_ADDR_ASCAST:.*]] = addrspacecast ptr %[[T_ADDR]] to ptr addrspace(4) -// OGCG-NEXT: store ptr %[[T]], ptr addrspace(4) %[[T_ADDR_ASCAST]], align 8 -// OGCG-NEXT: ret void +// LLVM-LABEL: define {{.*}}spir_func void @_Z4tmplIPU3AS0iEvT_( +// LLVM-SAME: ptr noundef %[[T:.*]]) +// LLVM: %[[T_ADDR:.*]] = alloca ptr, align 8 +// LLVM-NEXT: %[[T_ADDR_ASCAST:.*]] = addrspacecast ptr %[[T_ADDR]] to ptr addrspace(4) +// LLVM-NEXT: store ptr %[[T]], ptr addrspace(4) %[[T_ADDR_ASCAST]], align 8 +// LLVM-NEXT: ret void -// OGCG-LABEL: define {{.*}}spir_func void @_Z4tmplIPiEvT_( -// OGCG-SAME: ptr addrspace(4) noundef %[[T:.*]]) -// OGCG: %[[T_ADDR:.*]] = alloca ptr addrspace(4), align 8 -// OGCG-NEXT: %[[T_ADDR_ASCAST:.*]] = addrspacecast ptr %[[T_ADDR]] to ptr addrspace(4) -// OGCG-NEXT: store ptr addrspace(4) %[[T]], ptr addrspace(4) %[[T_ADDR_ASCAST]], align 8 -// OGCG-NEXT: ret void +// LLVM-LABEL: define {{.*}}spir_func void @_Z4tmplIPiEvT_( +// LLVM-SAME: ptr addrspace(4) noundef %[[T:.*]]) +// LLVM: %[[T_ADDR:.*]] = alloca ptr addrspace(4), align 8 +// LLVM-NEXT: %[[T_ADDR_ASCAST:.*]] = addrspacecast ptr %[[T_ADDR]] to ptr addrspace(4) +// LLVM-NEXT: store ptr addrspace(4) %[[T]], ptr addrspace(4) %[[T_ADDR_ASCAST]], align 8 +// LLVM-NEXT: ret void >From f45a0f0e0ef04e672f748e59e0b7c7357d9f13f1 Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris <[email protected]> Date: Wed, 7 Oct 2026 11:34:29 -0700 Subject: [PATCH 3/4] Fix nit Co-Authored-By: Claude Opus 5.5 (1M context) <[email protected]> --- clang/lib/CIR/CodeGen/CIRGenExprAggregate.cpp | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenExprAggregate.cpp b/clang/lib/CIR/CodeGen/CIRGenExprAggregate.cpp index 9e091568d2a7f..9006d9351aa8c 100644 --- a/clang/lib/CIR/CodeGen/CIRGenExprAggregate.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenExprAggregate.cpp @@ -751,10 +751,8 @@ void AggExprEmitter::emitArrayInit(Address destPtr, cir::ArrayType arrayTy, const QualType elementPtrType = cgf.getContext().getPointerType(elementType); const mlir::Type cirElementType = cgf.convertType(elementType); - const cir::PointerType cirElementPtrType = builder.getPointerTo( - cirElementType, - mlir::cast<cir::PointerType>(destPtr.getPointer().getType()) - .getAddrSpace()); + const cir::PointerType cirElementPtrType = + builder.getPointerTo(cirElementType, destPtr.getAddressSpace()); auto begin = cir::CastOp::create(builder, loc, cirElementPtrType, cir::CastKind::array_to_ptrdecay, >From 03cd6f7120896f3c076a8de21afb9310ad2ddff1 Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris <[email protected]> Date: Wed, 7 Oct 2026 14:51:06 -0700 Subject: [PATCH 4/4] [CIR][SYCL] Update null pointer checks after #228559 #228559 emits null pointers to named address spaces as an addrspacecast of the generic null pointer, matching classic CodeGen. Update the LLVM-CIR checks and drop the stale TODO. Co-Authored-By: Claude Opus 5.5 (1M context) <[email protected]> --- clang/test/CIR/CodeGenSYCL/address-space-conversions.cpp | 8 +++----- 1 file changed, 3 insertions(+), 5 deletions(-) diff --git a/clang/test/CIR/CodeGenSYCL/address-space-conversions.cpp b/clang/test/CIR/CodeGenSYCL/address-space-conversions.cpp index 5a2213a8809d2..b2a5468adaa61 100644 --- a/clang/test/CIR/CodeGenSYCL/address-space-conversions.cpp +++ b/clang/test/CIR/CodeGenSYCL/address-space-conversions.cpp @@ -9,9 +9,7 @@ // variables and parameters are allocated in the private address space and // accessed through the generic address space. // -// TODO(cir): Calls are missing the spir_func calling convention and null -// pointers to named address spaces are not emitted as an addrspacecast of the -// generic null pointer. +// TODO(cir): Calls are missing the spir_func calling convention. void bar(int &Data) {} void bar2(int &Data) {} @@ -122,8 +120,8 @@ void tmpl(T t) {} // LLVM-CIR-NEXT: %[[LOC_ASCAST:.*]] = addrspacecast ptr %[[LOC]] to ptr addrspace(4) // LLVM-CIR-NEXT: %[[GLOB_ASCAST:.*]] = addrspacecast ptr %[[GLOB]] to ptr addrspace(4) // LLVM-CIR-NEXT: %[[NOAS_ASCAST:.*]] = addrspacecast ptr %[[NOAS]] to ptr addrspace(4) -// LLVM-CIR-NEXT: store ptr addrspace(3) null, ptr addrspace(4) %[[LOC_ASCAST]], align 8 -// LLVM-CIR-NEXT: store ptr addrspace(1) null, ptr addrspace(4) %[[GLOB_ASCAST]], align 8 +// LLVM-CIR-NEXT: store ptr addrspace(3) addrspacecast (ptr addrspace(4) null to ptr addrspace(3)), ptr addrspace(4) %[[LOC_ASCAST]], align 8 +// LLVM-CIR-NEXT: store ptr addrspace(1) addrspacecast (ptr addrspace(4) null to ptr addrspace(1)), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 // LLVM-CIR-NEXT: %[[TMP0:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[GLOB_ASCAST]], align 8 // LLVM-CIR-NEXT: %[[TMP1:.*]] = addrspacecast ptr addrspace(1) %[[TMP0]] to ptr addrspace(4) // LLVM-CIR-NEXT: store ptr addrspace(4) %[[TMP1]], ptr addrspace(4) %[[NOAS_ASCAST]], align 8 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
