[clang] [CIR] Keep the address space in array ctor/dtor and new[]/delete[] (PR #231367)
Arseniy Obolenskiy via cfe-commits
cfe-commits at lists.llvm.org
Mon Oct 12 22:22:32 PDT 2026
https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/231367
>From d6abfdddd63cc4231f29fef47d4ba171bba314f9 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Tue, 13 Oct 2026 06:55:27 +0200
Subject: [PATCH] [CIR] Keep the address space in array ctor/dtor and
new[]/delete[]
Pointers to array elements lost the address space of the array, whichi broke the CIR verifier on spirv64-amd-amdhsa
---
clang/lib/CIR/CodeGen/CIRGenClass.cpp | 3 +-
clang/lib/CIR/CodeGen/CIRGenDecl.cpp | 8 +-
clang/lib/CIR/CodeGen/CIRGenItaniumCXXABI.cpp | 6 +-
clang/lib/CIR/Dialect/IR/CIRDialect.cpp | 4 +
.../CIR/Dialect/Transforms/CXXABILowering.cpp | 11 +-
.../Transforms/TargetLowering/CIRCXXABI.cpp | 6 +-
.../TargetLowering/LowerItaniumCXXABI.cpp | 6 +-
.../CodeGenHIP/array-structor-addrspace.hip | 109 ++++++++++++++++++
clang/test/CIR/IR/invalid-array-structor.cir | 30 +++++
9 files changed, 165 insertions(+), 18 deletions(-)
create mode 100644 clang/test/CIR/CodeGenHIP/array-structor-addrspace.hip
diff --git a/clang/lib/CIR/CodeGen/CIRGenClass.cpp b/clang/lib/CIR/CodeGen/CIRGenClass.cpp
index 679fe082ac37e..1af2357ed034a 100644
--- a/clang/lib/CIR/CodeGen/CIRGenClass.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenClass.cpp
@@ -724,7 +724,8 @@ void CIRGenFunction::emitCXXAggrConstructorCall(
mlir::Type elementType = arrayBase.getElementType();
while (auto maybeArrayTy = mlir::dyn_cast<cir::ArrayType>(elementType))
elementType = maybeArrayTy.getElementType();
- cir::PointerType ptrToElmType = builder.getPointerTo(elementType);
+ cir::PointerType ptrToElmType =
+ builder.getPointerTo(elementType, arrayBase.getAddressSpace());
bool useDynamicArrayCtor = true;
uint64_t constElementCount = 0;
diff --git a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
index f841cecbe0a85..87ac863c10aad 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
@@ -1117,7 +1117,9 @@ struct IrregularPartialArrayDestroy final : EHScopeStack::Cleanup {
QualType baseElementType = cgf.getContext().getBaseElementType(elementType);
mlir::Type cirElementType = cgf.convertTypeForMem(baseElementType);
- cir::PointerType ptrToElmType = builder.getPointerTo(cirElementType);
+ cir::PointerType ptrToElmType = builder.getPointerTo(
+ cirElementType,
+ mlir::cast<cir::PointerType>(arrayBegin.getType()).getAddrSpace());
mlir::Value begin = arrayBegin;
if (baseElementType != elementType) {
@@ -1270,7 +1272,9 @@ void CIRGenFunction::emitArrayDestroy(mlir::Value begin,
// Differently from LLVM traditional codegen, use a higher level
// representation instead of lowering directly to a loop.
mlir::Type cirElementType = convertTypeForMem(elementType);
- cir::PointerType ptrToElmType = builder.getPointerTo(cirElementType);
+ cir::PointerType ptrToElmType = builder.getPointerTo(
+ cirElementType,
+ mlir::cast<cir::PointerType>(begin.getType()).getAddrSpace());
auto regionBuilder = [&](mlir::OpBuilder &b, mlir::Location loc) {
mlir::BlockArgument arg =
diff --git a/clang/lib/CIR/CodeGen/CIRGenItaniumCXXABI.cpp b/clang/lib/CIR/CodeGen/CIRGenItaniumCXXABI.cpp
index 25526f0668f00..73461fa43752d 100644
--- a/clang/lib/CIR/CodeGen/CIRGenItaniumCXXABI.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenItaniumCXXABI.cpp
@@ -2228,9 +2228,6 @@ Address CIRGenItaniumCXXABI::initializeArrayCookie(CIRGenFunction &cgf,
QualType elementType) {
assert(requiresArrayCookie(e));
- // TODO: When sanitizer support is implemented, we'll need to
- // get the address space from `newPtr`.
- assert(!cir::MissingFeatures::addressSpace());
assert(!cir::MissingFeatures::sanitizers());
ASTContext &ctx = cgm.getASTContext();
@@ -2243,9 +2240,8 @@ Address CIRGenItaniumCXXABI::initializeArrayCookie(CIRGenFunction &cgf,
assert(cookieSize == getArrayCookieSizeImpl(elementType));
mlir::Type u8Ty = cgf.getBuilder().getUInt8Ty();
- cir::PointerType u8PtrTy = cgf.getBuilder().getUInt8PtrTy();
mlir::Value baseBytePtr =
- cgf.getBuilder().createBitcast(newPtr.getPointer(), u8PtrTy);
+ cgf.getBuilder().createPtrBitcast(newPtr.getPointer(), u8Ty);
// Compute an offset to the cookie.
CharUnits cookieOffset = cookieSize - sizeSize;
diff --git a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp
index 0fb9f585d2228..cf3b354a5dc40 100644
--- a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp
+++ b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp
@@ -542,6 +542,10 @@ template <typename Op> static LogicalResult verifyArrayCtorDtor(Op op) {
return op.emitOpError(
"block argument pointee type must match the innermost array "
"element type");
+
+ if (expectedEltPtrTy.getAddrSpace() != ptrTy.getAddrSpace())
+ return op.emitOpError(
+ "block argument address space must match the 'addr' address space");
}
return success();
diff --git a/clang/lib/CIR/Dialect/Transforms/CXXABILowering.cpp b/clang/lib/CIR/Dialect/Transforms/CXXABILowering.cpp
index fd4e1863de5e5..263f72ec3fe32 100644
--- a/clang/lib/CIR/Dialect/Transforms/CXXABILowering.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/CXXABILowering.cpp
@@ -677,7 +677,7 @@ mlir::LogicalResult CIRDeleteArrayOpABILowering::matchAndRewrite(
// total-size computation), so it must dominate both regions.
mlir::Value deletePtr;
mlir::Value numElements;
- cir::PointerType ptrTy;
+ auto ptrTy = mlir::cast<cir::PointerType>(loweredAddress.getType());
clang::CharUnits cookieSize;
mlir::DataLayout dl(op->getParentOfType<mlir::ModuleOp>());
unsigned ptrWidth =
@@ -685,14 +685,14 @@ mlir::LogicalResult CIRDeleteArrayOpABILowering::matchAndRewrite(
cir::IntType sizeTy = cirBuilder.getUIntNTy(ptrWidth);
if (cookieRequired) {
- ptrTy = mlir::cast<cir::PointerType>(loweredAddress.getType());
clang::CharUnits elementAlign =
clang::CharUnits::fromQuantity(op.getElementAlign());
cxxABI.readArrayCookie(loc, loweredAddress, elementAlign, dl, cirBuilder,
numElements, deletePtr, cookieSize);
} else {
- deletePtr = cir::CastOp::create(rewriter, loc, cirBuilder.getVoidPtrTy(),
- cir::CastKind::bitcast, loweredAddress);
+ deletePtr = cir::CastOp::create(
+ rewriter, loc, cirBuilder.getVoidPtrTy(ptrTy.getAddrSpace()),
+ cir::CastKind::bitcast, loweredAddress);
}
// Create a cleanup scope to wrap the ArrayDtor operation (if needed) and
@@ -719,12 +719,11 @@ mlir::LogicalResult CIRDeleteArrayOpABILowering::matchAndRewrite(
/*bodyBuilder=*/
[&](mlir::OpBuilder &b, mlir::Location l) {
if (dtorFn) {
- auto eltPtrTy = cir::PointerType::get(ptrTy.getPointee());
auto arrayDtor = cir::ArrayDtor::create(
b, l, loweredAddress, numElements,
[&](mlir::OpBuilder &bb, mlir::Location ll) {
mlir::Value arg =
- bb.getInsertionBlock()->addArgument(eltPtrTy, ll);
+ bb.getInsertionBlock()->addArgument(ptrTy, ll);
auto dtorCall = cir::CallOp::create(
bb, ll, dtorFn, cir::VoidType(), mlir::ValueRange{arg});
dtorCall.setCallingConv(getCalleeCallingConv(dtorFn));
diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/CIRCXXABI.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/CIRCXXABI.cpp
index df8ab49c067ec..a84da56724277 100644
--- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/CIRCXXABI.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/CIRCXXABI.cpp
@@ -28,9 +28,11 @@ void CIRCXXABI::readArrayCookie(mlir::Location loc, mlir::Value elementPtr,
CIRBaseBuilderTy &builder,
mlir::Value &numElements, mlir::Value &allocPtr,
clang::CharUnits &cookieSize) const {
- auto u8PtrTy = builder.getPointerTo(builder.getUIntNTy(8));
+ mlir::ptr::MemorySpaceAttrInterface addrSpace =
+ mlir::cast<cir::PointerType>(elementPtr.getType()).getAddrSpace();
+ auto u8PtrTy = builder.getPointerTo(builder.getUIntNTy(8), addrSpace);
auto ptrDiffTy = builder.getSIntNTy(getPtrSizeInBits());
- auto voidPtrTy = builder.getVoidPtrTy();
+ auto voidPtrTy = builder.getVoidPtrTy(addrSpace);
cookieSize = getArrayCookieSizeImpl(elementAlign);
diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerItaniumCXXABI.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerItaniumCXXABI.cpp
index aef66b0468490..bf46efca878fb 100644
--- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerItaniumCXXABI.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerItaniumCXXABI.cpp
@@ -906,7 +906,9 @@ mlir::Value LowerItaniumCXXABI::readArrayCookieImpl(
clang::CharUnits cookieAlignment, const mlir::DataLayout &dataLayout,
CIRBaseBuilderTy &builder) const {
unsigned ptrSizeInBits = getPtrSizeInBits();
- auto u8PtrTy = builder.getPointerTo(builder.getUIntNTy(8));
+ mlir::ptr::MemorySpaceAttrInterface addrSpace =
+ mlir::cast<cir::PointerType>(allocPtr.getType()).getAddrSpace();
+ auto u8PtrTy = builder.getPointerTo(builder.getUIntNTy(8), addrSpace);
auto ptrDiffTy = builder.getSIntNTy(ptrSizeInBits);
auto sizeTy = builder.getUIntNTy(ptrSizeInBits);
@@ -925,7 +927,7 @@ mlir::Value LowerItaniumCXXABI::readArrayCookieImpl(
countAlignment = cookieAlignment.alignmentAtOffset(countOffset);
}
- auto countPtrTy = cir::PointerType::get(sizeTy);
+ auto countPtrTy = cir::PointerType::get(sizeTy, addrSpace);
mlir::Value countPtr = cir::CastOp::create(
builder, loc, countPtrTy, cir::CastKind::bitcast, countBytePtr);
return cir::LoadOp::create(
diff --git a/clang/test/CIR/CodeGenHIP/array-structor-addrspace.hip b/clang/test/CIR/CodeGenHIP/array-structor-addrspace.hip
new file mode 100644
index 0000000000000..c5f9e41a1da7e
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/array-structor-addrspace.hip
@@ -0,0 +1,109 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fcuda-is-device -fclangir -emit-cir %s -o %t.cir
+// RUN: FileCheck --input-file=%t.cir %s -check-prefix=CIR
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fcuda-is-device -fclangir -emit-llvm %s -o %t-cir.ll
+// RUN: FileCheck --input-file=%t-cir.ll %s -check-prefix=LLVM
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fcuda-is-device -emit-llvm %s -o %t.ll
+// RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG
+
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fcuda-is-device -fclangir -fcxx-exceptions -fexceptions -DEXCEPTIONS -emit-cir %s -o %t-eh.cir
+// RUN: FileCheck --input-file=%t-eh.cir %s -check-prefix=CIR-EH
+
+// Array elements and array cookies are in the address space of the array.
+
+#define __device__ __attribute__((device))
+
+typedef __SIZE_TYPE__ size_t;
+__device__ void *operator new[](size_t);
+__device__ void operator delete[](void *);
+
+struct S {
+ __device__ S();
+ __device__ S(int);
+ __device__ ~S();
+ int x;
+};
+
+#ifndef EXCEPTIONS
+
+__device__ void local_array() { S arr[4]; }
+
+// CIR-LABEL: cir.func {{.*}}@_Z11local_arrayv
+// CIR: %[[ARR:.*]] = cir.cast address_space %{{.*}} : !cir.ptr<!cir.array<!rec_S x 4>> -> !cir.ptr<!cir.array<!rec_S x 4>, target_address_space(4)>
+// CIR: cir.cast array_to_ptrdecay %[[ARR]] : !cir.ptr<!cir.array<!rec_S x 4>, target_address_space(4)> -> !cir.ptr<!rec_S, target_address_space(4)>
+// CIR: cir.call @_ZN1SC1Ev(%{{.*}}) {{.*}} : (!cir.ptr<!rec_S, target_address_space(4)> {{.*}}) -> ()
+// CIR: cir.cast array_to_ptrdecay %[[ARR]] : !cir.ptr<!cir.array<!rec_S x 4>, target_address_space(4)> -> !cir.ptr<!rec_S, target_address_space(4)>
+// CIR: cir.call @_ZN1SD1Ev(%{{.*}}) {{.*}} : (!cir.ptr<!rec_S, target_address_space(4)> {{.*}}) -> ()
+
+// LLVM-LABEL: define {{.*}}@_Z11local_arrayv
+// LLVM: %[[ARR:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(4)
+// LLVM: getelementptr %struct.S, ptr addrspace(4) %[[ARR]], i32 0
+// LLVM: call spir_func addrspace(4) void @_ZN1SC1Ev(ptr addrspace(4) {{.*}})
+// LLVM: getelementptr %struct.S, ptr addrspace(4) %[[ARR]], i32 0
+// LLVM: call spir_func addrspace(4) void @_ZN1SD1Ev(ptr addrspace(4) {{.*}})
+
+// OGCG-LABEL: define {{.*}}@_Z11local_arrayv
+// OGCG: %[[ARR:.*]] = addrspacecast ptr %{{.*}} to ptr addrspace(4)
+// OGCG: getelementptr inbounds [4 x %struct.S], ptr addrspace(4) %[[ARR]], i32 0, i32 0
+// OGCG: call spir_func addrspace(4) void @_ZN1SC1Ev(ptr addrspace(4) {{.*}})
+// OGCG: getelementptr inbounds [4 x %struct.S], ptr addrspace(4) %[[ARR]], i32 0, i32 0
+// OGCG: call spir_func addrspace(4) void @_ZN1SD1Ev(ptr addrspace(4) {{.*}})
+
+__device__ S *new_array(int n) { return new S[n]; }
+
+// CIR-LABEL: cir.func {{.*}}@_Z9new_arrayi
+// CIR: %[[ALLOC:.*]] = cir.call @_Znam(%{{.*}}) {{.*}} -> (!cir.ptr<!void, target_address_space(4)> {{.*}})
+// CIR: %[[BYTES:.*]] = cir.cast bitcast %[[ALLOC]] : !cir.ptr<!void, target_address_space(4)> -> !cir.ptr<!u8i, target_address_space(4)>
+// CIR: %[[COOKIE:.*]] = cir.cast bitcast %[[BYTES]] : !cir.ptr<!u8i, target_address_space(4)> -> !cir.ptr<!u64i, target_address_space(4)>
+// CIR: cir.store align(8) %{{.*}}, %[[COOKIE]] : !u64i, !cir.ptr<!u64i, target_address_space(4)>
+// CIR: cir.ptr_stride %[[BYTES]], %{{.*}} : (!cir.ptr<!u8i, target_address_space(4)>, !s32i) -> !cir.ptr<!u8i, target_address_space(4)>
+// CIR: cir.call @_ZN1SC1Ev(%{{.*}}) {{.*}} : (!cir.ptr<!rec_S, target_address_space(4)> {{.*}}) -> ()
+
+// LLVM-LABEL: define {{.*}}@_Z9new_arrayi
+// LLVM: %[[ALLOC:.*]] = call spir_func {{.*}} ptr addrspace(4) @_Znam(
+// LLVM: store i64 %{{.*}}, ptr addrspace(4) %[[ALLOC]], align 8
+// LLVM: getelementptr i8, ptr addrspace(4) %[[ALLOC]], i64 8
+// LLVM: call spir_func addrspace(4) void @_ZN1SC1Ev(ptr addrspace(4) {{.*}})
+
+// OGCG-LABEL: define {{.*}}@_Z9new_arrayi
+// OGCG: %[[ALLOC:.*]] = call spir_func {{.*}} ptr addrspace(4) @_Znam(
+// OGCG: store i64 %{{.*}}, ptr addrspace(4) %[[ALLOC]], align 8
+// OGCG: getelementptr inbounds i8, ptr addrspace(4) %[[ALLOC]], i64 8
+// OGCG: call spir_func addrspace(4) void @_ZN1SC1Ev(ptr addrspace(4) {{.*}})
+
+__device__ void delete_array(S *p) { delete[] p; }
+
+// CIR-LABEL: cir.func {{.*}}@_Z12delete_arrayP1S
+// CIR: %[[BYTES:.*]] = cir.cast bitcast %{{.*}} : !cir.ptr<!rec_S, target_address_space(4)> -> !cir.ptr<!u8i, target_address_space(4)>
+// CIR: %[[ALLOC:.*]] = cir.ptr_stride %[[BYTES]], %{{.*}} : (!cir.ptr<!u8i, target_address_space(4)>, !s64i) -> !cir.ptr<!u8i, target_address_space(4)>
+// CIR: %[[VOID:.*]] = cir.cast bitcast %[[ALLOC]] : !cir.ptr<!u8i, target_address_space(4)> -> !cir.ptr<!void, target_address_space(4)>
+// CIR: %[[COOKIE:.*]] = cir.cast bitcast %[[ALLOC]] : !cir.ptr<!u8i, target_address_space(4)> -> !cir.ptr<!u64i, target_address_space(4)>
+// CIR: cir.load align(4) %[[COOKIE]] : !cir.ptr<!u64i, target_address_space(4)>, !u64i
+// CIR: cir.call @_ZN1SD1Ev(%{{.*}}) {{.*}} : (!cir.ptr<!rec_S, target_address_space(4)>) -> ()
+// CIR: cir.call @_ZdaPvm(%[[VOID]], %{{.*}}) {{.*}} : (!cir.ptr<!void, target_address_space(4)>, !u64i) -> ()
+
+// LLVM-LABEL: define {{.*}}@_Z12delete_arrayP1S
+// LLVM: %[[ALLOC:.*]] = getelementptr i8, ptr addrspace(4) %{{.*}}, i64 -8
+// LLVM: load i64, ptr addrspace(4) %[[ALLOC]], align 4
+// LLVM: call spir_func addrspace(4) void @_ZN1SD1Ev(ptr addrspace(4) %{{.*}})
+// LLVM: call spir_func addrspace(4) void @_ZdaPvm(ptr addrspace(4) %[[ALLOC]], i64 %{{.*}})
+
+// OGCG-LABEL: define {{.*}}@_Z12delete_arrayP1S
+// OGCG: %[[ALLOC:.*]] = getelementptr inbounds i8, ptr addrspace(4) %{{.*}}, i64 -8
+// OGCG: load i64, ptr addrspace(4) %[[ALLOC]], align 4
+// OGCG: call spir_func addrspace(4) void @_ZN1SD1Ev(ptr addrspace(4) {{.*}})
+// OGCG: call spir_func addrspace(4) void @_ZdaPvm(ptr addrspace(4) noundef %[[ALLOC]], i64 noundef %{{.*}})
+
+#else
+
+__device__ void partial_init() { S arr[4] = {1}; }
+
+// CIR-EH-LABEL: cir.func {{.*}}@_Z12partial_initv
+// CIR-EH: %[[BEGIN:.*]] = cir.cast array_to_ptrdecay %{{.*}} : !cir.ptr<!cir.array<!rec_S x 4>, target_address_space(4)> -> !cir.ptr<!rec_S, target_address_space(4)>
+// CIR-EH: } cleanup eh {
+// CIR-EH: %[[END:.*]] = cir.load align(8) %{{.*}} : !cir.ptr<!cir.ptr<!rec_S, target_address_space(4)>, target_address_space(4)>, !cir.ptr<!rec_S, target_address_space(4)>
+// CIR-EH: cir.cmp ne %[[END]], %[[BEGIN]] : !cir.ptr<!rec_S, target_address_space(4)>
+// CIR-EH: cir.ptr_stride %{{.*}}, %{{.*}} : (!cir.ptr<!rec_S, target_address_space(4)>, !s64i) -> !cir.ptr<!rec_S, target_address_space(4)>
+// CIR-EH: cir.call @_ZN1SD1Ev(%{{.*}}) {{.*}} : (!cir.ptr<!rec_S, target_address_space(4)> {{.*}}) -> ()
+
+#endif
diff --git a/clang/test/CIR/IR/invalid-array-structor.cir b/clang/test/CIR/IR/invalid-array-structor.cir
index aad53dc5e5ed3..1a123978f2d2c 100644
--- a/clang/test/CIR/IR/invalid-array-structor.cir
+++ b/clang/test/CIR/IR/invalid-array-structor.cir
@@ -187,6 +187,36 @@ module {
// -----
+!u8i = !cir.int<u, 8>
+!rec_S = !cir.struct<"S" {pad !u8i}>
+
+module {
+ cir.func @bad_static_addrspace_mismatch_ctor(%p: !cir.ptr<!cir.array<!rec_S x 10>, target_address_space(1)>) {
+ // expected-error at +1 {{'cir.array.ctor' op block argument address space must match the 'addr' address space}}
+ cir.array.ctor %p : !cir.ptr<!cir.array<!rec_S x 10>, target_address_space(1)> {
+ ^bb0(%e: !cir.ptr<!rec_S>):
+ }
+ cir.return
+ }
+}
+
+// -----
+
+!u8i = !cir.int<u, 8>
+!rec_S = !cir.struct<"S" {pad !u8i}>
+
+module {
+ cir.func @bad_static_addrspace_mismatch_dtor(%p: !cir.ptr<!cir.array<!rec_S x 10>, target_address_space(1)>) {
+ // expected-error at +1 {{'cir.array.dtor' op block argument address space must match the 'addr' address space}}
+ cir.array.dtor %p : !cir.ptr<!cir.array<!rec_S x 10>, target_address_space(1)> {
+ ^bb0(%e: !cir.ptr<!rec_S>):
+ }
+ cir.return
+ }
+}
+
+// -----
+
!u8i = !cir.int<u, 8>
!u64i = !cir.int<u, 64>
!rec_S = !cir.struct<"S" {pad !u8i}>
More information about the cfe-commits
mailing list