[clang] Cir/sycl local generic addrspace (PR #229281)

[Konstantinos Parasyris via cfe-commits]

https://github.com/koparasy updated https://github.com/llvm/llvm-project/pull/229281

>From f464adb5e42a5fb9a0c21f53f135ede6e2490677 Mon Sep 17 00:00:00 2001
From: Konstantinos Parasyris <konstantinos.parasyris at intel.com>
Date: Sun, 4 Oct 2026 12:16:37 -0700
Subject: [PATCH 1/3] [CIR] Lower language address spaces nested in record
 members

TargetLowering now derives from RecordRewritingTypeConverter, so it rebuilds
records whose members reach a language address space and keeps the rest. This
fixes "member type mismatch" for OpenCL and SYCL records with address-space
qualified pointer members.
---
 .../Transforms/RecordTypeConverter.cpp        |  50 ++++++
 .../Dialect/Transforms/RecordTypeConverter.h  |  15 ++
 .../CIR/Dialect/Transforms/TargetLowering.cpp | 162 +++++++++---------
 .../CIR/CodeGenOpenCL/record-address-space.cl |  77 +++++++++
 .../CIR/CodeGenSYCL/record-address-space.cpp  |  34 ++++
 5 files changed, 258 insertions(+), 80 deletions(-)
 create mode 100644 clang/test/CIR/CodeGenOpenCL/record-address-space.cl
 create mode 100644 clang/test/CIR/CodeGenSYCL/record-address-space.cpp

diff --git a/clang/lib/CIR/Dialect/Transforms/RecordTypeConverter.cpp b/clang/lib/CIR/Dialect/Transforms/RecordTypeConverter.cpp
index 52c355c340519bb..8f14bada8957c7b 100644
--- a/clang/lib/CIR/Dialect/Transforms/RecordTypeConverter.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/RecordTypeConverter.cpp
@@ -108,8 +108,58 @@ RecordRewritingTypeConverter::convertRecordMemberTypes(cir::RecordType type) {
   return loweredMemberTypes;
 }
 
+bool RecordRewritingTypeConverter::reachesChangingType(
+    mlir::Type type, llvm::SmallPtrSetImpl<mlir::Type> &visiting) const {
+  if (typeMayChange(type))
+    return true;
+
+  // The sub-element walk skips record members. Stop on cycles.
+  if (auto rt = mlir::dyn_cast<cir::RecordType>(type)) {
+    if (!visiting.insert(type).second)
+      return false;
+    for (mlir::Type member : rt.getMembers())
+      if (reachesChangingType(member, visiting))
+        return true;
+    if (auto u = mlir::dyn_cast<cir::UnionType>(type))
+      if (mlir::Type pad = u.getPadding())
+        return reachesChangingType(pad, visiting);
+    return false;
+  }
+
+  bool found = false;
+  type.walkImmediateSubElements([](mlir::Attribute) {},
+                                [&](mlir::Type sub) {
+                                  found = found ||
+                                          reachesChangingType(sub, visiting);
+                                });
+  return found;
+}
+
+bool RecordRewritingTypeConverter::recordNeedsConversion(cir::RecordType type) {
+  {
+    std::shared_lock<decltype(recordNeedsConversionMutex)> lock(
+        recordNeedsConversionMutex);
+    auto it = recordNeedsConversionCache.find(type);
+    if (it != recordNeedsConversionCache.end())
+      return it->second;
+  }
+
+  // Nested results may be partial inside a cycle, so cache only this one.
+  llvm::SmallPtrSet<mlir::Type, 8> visiting;
+  bool result = reachesChangingType(type, visiting);
+
+  std::unique_lock<decltype(recordNeedsConversionMutex)> lock(
+      recordNeedsConversionMutex);
+  recordNeedsConversionCache[type] = result;
+  return result;
+}
+
 cir::RecordType
 RecordRewritingTypeConverter::convertRecordType(cir::RecordType type) {
+  // Keep records that reach no changing type, including through nested ones.
+  if (!recordNeedsConversion(type))
+    return type;
+
   // Unnamed record types can't be referred to recursively, so we can just
   // convert this one. It also doesn't have uniqueness problems, so we can
   // just do a conversion on it.
diff --git a/clang/lib/CIR/Dialect/Transforms/RecordTypeConverter.h b/clang/lib/CIR/Dialect/Transforms/RecordTypeConverter.h
index 1b7b1fbc1ad34c0..569a96a604e337e 100644
--- a/clang/lib/CIR/Dialect/Transforms/RecordTypeConverter.h
+++ b/clang/lib/CIR/Dialect/Transforms/RecordTypeConverter.h
@@ -17,6 +17,7 @@
 #include "mlir/Transforms/DialectConversion.h"
 #include "clang/CIR/Dialect/IR/CIRTypes.h"
 #include "llvm/ADT/DenseMap.h"
+#include "llvm/ADT/SmallPtrSet.h"
 #include "llvm/ADT/SmallVector.h"
 #include "llvm/Support/RWMutex.h"
 
@@ -36,6 +37,8 @@ namespace cir {
 /// Subclasses add conversions for the types they rewrite. The type converter
 /// tries the most recently added conversion first, so a subclass can also
 /// replace one of the conversions registered here.
+///
+/// Every record is rebuilt unless typeMayChange() is overridden.
 class RecordRewritingTypeConverter : public mlir::TypeConverter {
 public:
   explicit RecordRewritingTypeConverter(mlir::MLIRContext &context);
@@ -44,7 +47,15 @@ class RecordRewritingTypeConverter : public mlir::TypeConverter {
   /// Remove the temporary name of every record rebuilt by this converter.
   void restoreRecordTypeNames();
 
+protected:
+  /// Whether \p type itself, ignoring nested types, may change. Records are
+  /// rebuilt only if some type they reach returns true.
+  virtual bool typeMayChange(mlir::Type type) const { return true; }
+
 private:
+  bool recordNeedsConversion(cir::RecordType type);
+  bool reachesChangingType(mlir::Type type,
+                           llvm::SmallPtrSetImpl<mlir::Type> &visiting) const;
   cir::RecordType convertRecordType(cir::RecordType type);
   llvm::SmallVector<mlir::Type> convertRecordMemberTypes(cir::RecordType type);
   llvm::SmallVector<cir::RecordType> &getCurrentThreadRecursiveStack();
@@ -64,6 +75,10 @@ class RecordRewritingTypeConverter : public mlir::TypeConverter {
   // functions.
   llvm::SmallVector<cir::RecordType> convertedRecordTypes;
   llvm::sys::SmartRWMutex<true> recordTypeMutex;
+
+  // Cached results of recordNeedsConversion().
+  llvm::DenseMap<mlir::Type, bool> recordNeedsConversionCache;
+  llvm::sys::SmartRWMutex<true> recordNeedsConversionMutex;
 };
 
 } // namespace cir
diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering.cpp
index 769df5572f285e5..76ca7114ee5ceae 100644
--- a/clang/lib/CIR/Dialect/Transforms/TargetLowering.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering.cpp
@@ -10,9 +10,11 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include "RecordTypeConverter.h"
 #include "TargetLowering/LowerModule.h"
 #include "TargetLowering/TargetLoweringInfo.h"
 
+#include "mlir/IR/AttrTypeSubElements.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/Support/LLVM.h"
 #include "mlir/Transforms/DialectConversion.h"
@@ -38,13 +40,72 @@ struct TargetLoweringPass
   void runOnOperation() override;
 };
 
+/// Converts LangAddressSpaceAttr → TargetAddressSpaceAttr inside types,
+/// rebuilding only records that reach a language address space.
+class TargetLoweringTypeConverter : public cir::RecordRewritingTypeConverter {
+public:
+  TargetLoweringTypeConverter(mlir::MLIRContext &ctx,
+                              const cir::TargetLoweringInfo &targetInfo)
+      : RecordRewritingTypeConverter(ctx), targetInfo(targetInfo) {
+    addConversion([this](cir::PointerType type) -> mlir::Type {
+      mlir::Type pointee = convertType(type.getPointee());
+      if (!pointee)
+        return {};
+      mlir::ptr::MemorySpaceAttrInterface addrSpace = type.getAddrSpace();
+      if (auto langAS =
+              mlir::dyn_cast_if_present<cir::LangAddressSpaceAttr>(addrSpace))
+        addrSpace = lowerAddrSpace(type.getContext(), langAS);
+      return cir::PointerType::get(type.getContext(), pointee, addrSpace);
+    });
+  }
+
+  mlir::ptr::MemorySpaceAttrInterface
+  lowerAddrSpace(mlir::MLIRContext *ctx,
+                 cir::LangAddressSpaceAttr langAS) const {
+    unsigned targetAS =
+        targetInfo.getTargetAddrSpaceFromCIRAddrSpace(langAS.getValue());
+    if (targetAS == 0)
+      return {};
+    return cir::TargetAddressSpaceAttr::get(ctx, targetAS);
+  }
+
+  /// Convert the types nested in \p attr, e.g. the type of a constant.
+  mlir::Attribute convertAttr(mlir::Attribute attr) const {
+    return attrReplacer.replace(attr);
+  }
+
+protected:
+  bool typeMayChange(mlir::Type type) const override {
+    auto ptrTy = mlir::dyn_cast<cir::PointerType>(type);
+    return ptrTy && mlir::isa_and_present<cir::LangAddressSpaceAttr>(
+                        ptrTy.getAddrSpace());
+  }
+
+private:
+  const cir::TargetLoweringInfo &targetInfo;
+  mutable mlir::AttrTypeReplacer attrReplacer = [this] {
+    mlir::AttrTypeReplacer replacer;
+    replacer.addReplacement(
+        [this](mlir::Type type)
+            -> std::optional<std::pair<mlir::Type, mlir::WalkResult>> {
+          mlir::Type converted = convertType(type);
+          if (!converted)
+            return std::nullopt;
+          // Already converted; don't walk into it.
+          return std::make_pair(converted, mlir::WalkResult::skip());
+        });
+    return replacer;
+  }();
+};
+
 /// A generic target lowering pattern that matches any CIR op whose operand or
 /// result types need address space conversion. Clones the op with converted
 /// types.
 class CIRGenericTargetLoweringPattern : public mlir::ConversionPattern {
 public:
-  CIRGenericTargetLoweringPattern(mlir::MLIRContext *context,
-                                  const mlir::TypeConverter &typeConverter)
+  CIRGenericTargetLoweringPattern(
+      mlir::MLIRContext *context,
+      const TargetLoweringTypeConverter &typeConverter)
       : mlir::ConversionPattern(typeConverter, MatchAnyOpTypeTag(),
                                 /*benefit=*/1, context) {}
 
@@ -73,9 +134,7 @@ class CIRGenericTargetLoweringPattern : public mlir::ConversionPattern {
     mlir::OperationState loweredOpState(op->getLoc(), op->getName());
     loweredOpState.addOperands(operands);
 
-    // Preserve auxiliary metadata verbatim. Convert only inherent TypeAttrs and
-    // pointer constants so address-space-bearing operation semantics (e.g.
-    // AllocaOp's allocaType) stay in sync with the converted result types.
+    // Preserve auxiliary metadata verbatim.
     loweredOpState.propertiesAttr = op->getPropertiesAsAttribute();
     loweredOpState.addAttributes(op->getDiscardableAttrDictionary().getValue());
 
@@ -96,15 +155,13 @@ class CIRGenericTargetLoweringPattern : public mlir::ConversionPattern {
     }
 
     mlir::Operation *loweredOp = rewriter.create(loweredOpState);
+    // Keep inherent attributes (e.g. allocaType, constants) in sync with the
+    // converted types.
+    const auto *tlTypeConverter =
+        static_cast<const TargetLoweringTypeConverter *>(typeConverter);
     loweredOp->getName().walkInherentAttrs(
         loweredOp, [&](llvm::StringRef, mlir::Attribute &attr) {
-          if (auto typeAttr = mlir::dyn_cast<mlir::TypeAttr>(attr))
-            attr = mlir::TypeAttr::get(
-                typeConverter->convertType(typeAttr.getValue()));
-          else if (auto ptrAttr = mlir::dyn_cast<cir::ConstPtrAttr>(attr))
-            attr = cir::ConstPtrAttr::get(
-                typeConverter->convertType(ptrAttr.getType()),
-                ptrAttr.getValue());
+          attr = tlTypeConverter->convertAttr(attr);
         });
     rewriter.replaceOp(op, loweredOp);
     return mlir::success();
@@ -116,34 +173,26 @@ class CIRGenericTargetLoweringPattern : public mlir::ConversionPattern {
 /// TypeConverter won't reach it automatically.
 class CIRGlobalOpTargetLowering
     : public mlir::OpConversionPattern<cir::GlobalOp> {
-  const cir::TargetLoweringInfo &targetInfo;
-
 public:
   CIRGlobalOpTargetLowering(mlir::MLIRContext *context,
-                            const mlir::TypeConverter &typeConverter,
-                            const cir::TargetLoweringInfo &targetInfo)
+                            const TargetLoweringTypeConverter &typeConverter)
       : mlir::OpConversionPattern<cir::GlobalOp>(typeConverter, context,
-                                                 /*benefit=*/1),
-        targetInfo(targetInfo) {}
+                                                 /*benefit=*/1) {}
 
   mlir::LogicalResult
   matchAndRewrite(cir::GlobalOp op, OpAdaptor adaptor,
                   mlir::ConversionPatternRewriter &rewriter) const override {
-    mlir::Type loweredSymTy = getTypeConverter()->convertType(op.getSymType());
+    const auto *typeConverter =
+        static_cast<const TargetLoweringTypeConverter *>(getTypeConverter());
+    mlir::Type loweredSymTy = typeConverter->convertType(op.getSymType());
     if (!loweredSymTy)
       return mlir::failure();
 
     // Convert the addr_space attribute.
     mlir::ptr::MemorySpaceAttrInterface addrSpace = op.getAddrSpaceAttr();
     if (auto langAS =
-            mlir::dyn_cast_if_present<cir::LangAddressSpaceAttr>(addrSpace)) {
-      unsigned targetAS =
-          targetInfo.getTargetAddrSpaceFromCIRAddrSpace(langAS.getValue());
-      addrSpace =
-          targetAS == 0
-              ? nullptr
-              : cir::TargetAddressSpaceAttr::get(op.getContext(), targetAS);
-    }
+            mlir::dyn_cast_if_present<cir::LangAddressSpaceAttr>(addrSpace))
+      addrSpace = typeConverter->lowerAddrSpace(op.getContext(), langAS);
 
     // Only rewrite if something actually changed.
     if (loweredSymTy == op.getSymType() && addrSpace == op.getAddrSpaceAttr())
@@ -152,6 +201,8 @@ class CIRGlobalOpTargetLowering
     auto newOp = mlir::cast<cir::GlobalOp>(rewriter.clone(*op.getOperation()));
     newOp.setSymType(loweredSymTy);
     newOp.setAddrSpaceAttr(addrSpace);
+    if (mlir::Attribute init = newOp.getInitialValueAttr())
+      newOp.setInitialValueAttr(typeConverter->convertAttr(init));
     rewriter.replaceOp(op, newOp);
     return mlir::success();
   }
@@ -218,55 +269,6 @@ static void convertSyncScopeIfPresent(mlir::Operation *op,
   }
 }
 
-/// Prepare the type converter for the target lowering pass.
-/// Converts LangAddressSpaceAttr → TargetAddressSpaceAttr inside pointer types.
-static void
-prepareTargetLoweringTypeConverter(mlir::TypeConverter &converter,
-                                   const cir::TargetLoweringInfo &targetInfo) {
-  converter.addConversion([](mlir::Type type) { return type; });
-
-  converter.addConversion([&converter,
-                           &targetInfo](cir::PointerType type) -> mlir::Type {
-    mlir::Type pointee = converter.convertType(type.getPointee());
-    if (!pointee)
-      return {};
-    auto addrSpace = type.getAddrSpace();
-    if (auto langAS =
-            mlir::dyn_cast_if_present<cir::LangAddressSpaceAttr>(addrSpace)) {
-      unsigned targetAS =
-          targetInfo.getTargetAddrSpaceFromCIRAddrSpace(langAS.getValue());
-      addrSpace =
-          targetAS == 0
-              ? nullptr
-              : cir::TargetAddressSpaceAttr::get(type.getContext(), targetAS);
-    }
-    return cir::PointerType::get(type.getContext(), pointee, addrSpace);
-  });
-
-  converter.addConversion([&converter](cir::ArrayType type) -> mlir::Type {
-    mlir::Type loweredElementType =
-        converter.convertType(type.getElementType());
-    if (!loweredElementType)
-      return {};
-    return cir::ArrayType::get(loweredElementType, type.getSize());
-  });
-
-  converter.addConversion([&converter](cir::FuncType type) -> mlir::Type {
-    llvm::SmallVector<mlir::Type> loweredInputTypes;
-    loweredInputTypes.reserve(type.getNumInputs());
-    if (mlir::failed(
-            converter.convertTypes(type.getInputs(), loweredInputTypes)))
-      return {};
-
-    mlir::Type loweredReturnType = converter.convertType(type.getReturnType());
-    if (!loweredReturnType)
-      return {};
-
-    return cir::FuncType::get(loweredInputTypes, loweredReturnType,
-                              /*isVarArg=*/type.getVarArg());
-  });
-}
-
 static void
 populateTargetLoweringConversionTarget(mlir::ConversionTarget &target,
                                        const mlir::TypeConverter &tc) {
@@ -311,12 +313,10 @@ void TargetLoweringPass::runOnOperation() {
   });
 
   // Address space conversion: LangAddressSpaceAttr → TargetAddressSpaceAttr.
-  mlir::TypeConverter typeConverter;
-  prepareTargetLoweringTypeConverter(typeConverter, targetInfo);
+  TargetLoweringTypeConverter typeConverter(*mod.getContext(), targetInfo);
 
   mlir::RewritePatternSet patterns(mod.getContext());
-  patterns.add<CIRGlobalOpTargetLowering>(mod.getContext(), typeConverter,
-                                          targetInfo);
+  patterns.add<CIRGlobalOpTargetLowering>(mod.getContext(), typeConverter);
   patterns.add<CIRFuncOpTargetLowering>(typeConverter, mod.getContext());
   patterns.add<CIRGenericTargetLoweringPattern>(mod.getContext(),
                                                 typeConverter);
@@ -330,6 +330,8 @@ void TargetLoweringPass::runOnOperation() {
 
   if (failed(mlir::applyPartialConversion(ops, target, std::move(patterns))))
     signalPassFailure();
+
+  typeConverter.restoreRecordTypeNames();
 }
 
 std::unique_ptr<Pass> mlir::createTargetLoweringPass() {
diff --git a/clang/test/CIR/CodeGenOpenCL/record-address-space.cl b/clang/test/CIR/CodeGenOpenCL/record-address-space.cl
new file mode 100644
index 000000000000000..8ce869a47a3b11a
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenCL/record-address-space.cl
@@ -0,0 +1,77 @@
+// RUN: %clang_cc1 -triple spirv64-unknown-unknown -cl-std=CL2.0 -fclangir \
+// RUN:   -emit-cir -mmlir --mlir-print-ir-before=cir-target-lowering %s \
+// RUN:   -o %t.cir 2> %t.pre.cir
+// RUN: FileCheck %s --check-prefix=PRE --input-file=%t.pre.cir
+// RUN: FileCheck %s --check-prefix=POST --input-file=%t.cir
+// RUN: %clang_cc1 -triple spirv64-unknown-unknown -cl-std=CL2.0 -fclangir \
+// RUN:   -emit-llvm -O0 %s -o %t-cir.ll
+// RUN: FileCheck %s --check-prefixes=LLVM,CIR-LLVM --input-file=%t-cir.ll
+// RUN: %clang_cc1 -triple spirv64-unknown-unknown -cl-std=CL2.0 \
+// RUN:   -emit-llvm -O0 %s -o %t.ll
+// RUN: FileCheck %s --check-prefixes=LLVM,OGCG --input-file=%t.ll
+
+// Language address spaces in record members are lowered; other records are
+// kept.
+
+struct S { __global int *p; int n; };
+struct Outer { struct S s; __local int *lp; };
+struct Node { __global struct Node *next; int v; };
+union U { __global int *p; long l; };
+struct Plain { int a; float b; };
+
+// PRE-DAG: !rec_S = !cir.struct<"S" {data !cir.ptr<!s32i, lang_address_space(offload_global)>, data !s32i}>
+// PRE-DAG: !rec_Outer = !cir.struct<"Outer" {data !rec_S, data !cir.ptr<!s32i, lang_address_space(offload_local)>}>
+// PRE-DAG: !rec_Node = !cir.struct<"Node" {data !cir.ptr<!cir.struct<"Node">, lang_address_space(offload_global)>, data !s32i}>
+// PRE-DAG: !rec_U = !cir.union<"U" {data !cir.ptr<!s32i, lang_address_space(offload_global)>, data !s64i}>
+// PRE-DAG: !rec_Plain = !cir.struct<"Plain" {data !s32i, data !cir.float}>
+
+// POST-DAG: !rec_S = !cir.struct<"S" {data !cir.ptr<!s32i, target_address_space(1)>, data !s32i}>
+// POST-DAG: !rec_Outer = !cir.struct<"Outer" {data !rec_S, data !cir.ptr<!s32i, target_address_space(3)>}>
+// POST-DAG: !rec_Node = !cir.struct<"Node" {data !cir.ptr<!cir.struct<"Node">, target_address_space(1)>, data !s32i}>
+// POST-DAG: !rec_U = !cir.union<"U" {data !cir.ptr<!s32i, target_address_space(1)>, data !s64i}>
+// POST-DAG: !rec_Plain = !cir.struct<"Plain" {data !s32i, data !cir.float}>
+
+// LLVM-DAG: %struct.S = type { ptr addrspace(1), i32 }
+// LLVM-DAG: %struct.Outer = type { %struct.S, ptr addrspace(3) }
+// LLVM-DAG: %struct.Node = type { ptr addrspace(1), i32 }
+// LLVM-DAG: %union.U = type { ptr addrspace(1) }
+// LLVM-DAG: %struct.Plain = type { i32, float }
+
+__global int gi;
+__global struct S gs = {&gi, 4};
+
+// The initializer type follows the record.
+
+// POST: cir.global external target_address_space(1) @gs = #cir.const_record<{#cir.global_view<@gi> : !cir.ptr<!s32i, target_address_space(1)>, #cir.int<4> : !s32i}> : !rec_S
+
+// CIR-LLVM: @gs = addrspace(1) global %struct.S { ptr addrspace(1) @gi, i32 4 }, align 8
+// OGCG: @gs = addrspace(1) global { ptr addrspace(1), i32, [4 x i8] } { ptr addrspace(1) @gi, i32 4, [4 x i8] zeroinitializer }, align 8
+
+kernel void k(__global int *g, __local int *l, __global struct Node *n) {
+  struct S s;
+  s.p = g;
+  struct Outer o;
+  o.lp = l;
+  struct Node node;
+  node.next = n;
+  union U u;
+  u.p = g;
+  struct Plain plain;
+  plain.a = 1;
+}
+
+// POST: cir.get_member %{{.*}}[0] {name = "p"} : !cir.ptr<!rec_S> -> !cir.ptr<!cir.ptr<!s32i, target_address_space(1)>>
+// POST: cir.get_member %{{.*}}[1] {name = "lp"} : !cir.ptr<!rec_Outer> -> !cir.ptr<!cir.ptr<!s32i, target_address_space(3)>>
+// POST: cir.get_member %{{.*}}[0] {name = "next"} : !cir.ptr<!rec_Node> -> !cir.ptr<!cir.ptr<!rec_Node, target_address_space(1)>>
+// POST: cir.get_member %{{.*}}[0] {name = "p"} : !cir.ptr<!rec_U> -> !cir.ptr<!cir.ptr<!s32i, target_address_space(1)>>
+// POST: cir.get_member %{{.*}}[0] {name = "a"} : !cir.ptr<!rec_Plain> -> !cir.ptr<!s32i>
+
+// LLVM: %[[P:.*]] = getelementptr inbounds nuw %struct.S, ptr %{{.*}}, i32 0, i32 0
+// LLVM: store ptr addrspace(1) %{{.*}}, ptr %[[P]], align 8
+// LLVM: %[[LP:.*]] = getelementptr inbounds nuw %struct.Outer, ptr %{{.*}}, i32 0, i32 1
+// LLVM: store ptr addrspace(3) %{{.*}}, ptr %[[LP]], align 8
+// LLVM: %[[NEXT:.*]] = getelementptr inbounds nuw %struct.Node, ptr %{{.*}}, i32 0, i32 0
+// LLVM: store ptr addrspace(1) %{{.*}}, ptr %[[NEXT]], align 8
+// LLVM: store ptr addrspace(1) %{{.*}}, ptr %{{.*}}, align 8
+// LLVM: %[[A:.*]] = getelementptr inbounds nuw %struct.Plain, ptr %{{.*}}, i32 0, i32 0
+// LLVM: store i32 1, ptr %[[A]], align 4
diff --git a/clang/test/CIR/CodeGenSYCL/record-address-space.cpp b/clang/test/CIR/CodeGenSYCL/record-address-space.cpp
new file mode 100644
index 000000000000000..292aab806d9424b
--- /dev/null
+++ b/clang/test/CIR/CodeGenSYCL/record-address-space.cpp
@@ -0,0 +1,34 @@
+// RUN: %clang_cc1 -std=c++20 -fsycl-is-device -triple spirv64-unknown-unknown \
+// RUN:   -fclangir -emit-cir -mmlir --mlir-print-ir-before=cir-target-lowering \
+// RUN:   %s -o %t.cir 2> %t.pre.cir
+// RUN: FileCheck %s --check-prefix=PRE --input-file=%t.pre.cir
+// RUN: FileCheck %s --check-prefix=POST --input-file=%t.cir
+// RUN: %clang_cc1 -std=c++20 -fsycl-is-device -triple spirv64-unknown-unknown \
+// RUN:   -fclangir -emit-llvm %s -o %t-cir.ll
+// RUN: FileCheck %s --check-prefix=LLVM --input-file=%t-cir.ll
+// RUN: %clang_cc1 -std=c++20 -fsycl-is-device -triple spirv64-unknown-unknown \
+// RUN:   -emit-llvm %s -o %t.ll
+// RUN: FileCheck %s --check-prefix=LLVM --input-file=%t.ll
+
+// Language address spaces in record members are lowered.
+
+struct S {
+  __attribute__((opencl_global)) int *p;
+  int n;
+};
+
+// PRE: !rec_S = !cir.struct<"S" {data !cir.ptr<!s32i, lang_address_space(offload_global)>, data !s32i}>
+// POST: !rec_S = !cir.struct<"S" {data !cir.ptr<!s32i, target_address_space(1)>, data !s32i}>
+// LLVM: %struct.S = type { ptr addrspace(1), i32 }
+
+[[clang::sycl_external]] void store_member(__attribute__((opencl_global)) int *g) {
+  S s;
+  s.p = g;
+}
+
+// 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)>>
+
+// LLVM-LABEL: define {{.*}}@_Z12store_memberPU3AS1i
+// LLVM: %[[P:.*]] = getelementptr inbounds nuw %struct.S, ptr {{.*}}, i32 0, i32 0
+// LLVM: store ptr addrspace(1) %{{.*}}, ptr {{.*}}%[[P]], align 8

>From 1212e7cea7acd4ca3cdaa2973b0a96ebe7b4c55a Mon Sep 17 00:00:00 2001
From: Konstantinos Parasyris <konstantinos.parasyris at intel.com>
Date: Mon, 5 Oct 2026 17:35:17 -0700
Subject: [PATCH 2/3] [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) <noreply at anthropic.com>
---
 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 1f738f9737d31fd..c9f5517341cd06c 100644
--- a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp
@@ -2951,6 +2951,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 81a6eae7effcf91..0c82d0bb93289ba 100644
--- a/clang/lib/CIR/CodeGen/CIRGenExprAggregate.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenExprAggregate.cpp
@@ -724,8 +724,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 382a0008113e225..5e50173058d222a 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 dd92cfab285f19a..eced98bad1210a8 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 000000000000000..4ea1e603cd87a97
--- /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 000000000000000..22032fcfa070e29
--- /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 000000000000000..0c4649d07517639
--- /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 292aab806d9424b..ae881857972a8a6 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 31daefbc85da8da8a72f0f175bf07981c36ae3b7 Mon Sep 17 00:00:00 2001
From: Konstantinos Parasyris <konstantinos.parasyris at intel.com>
Date: Mon, 5 Oct 2026 17:40:54 -0700
Subject: [PATCH 3/3] [CIR][SYCL] Share identical LLVM checks between CIR and
 classic CodeGen

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply at anthropic.com>
---
 .../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 4ea1e603cd87a97..5a2213a8809d26c 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