llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clangir Author: Arseniy Obolenskiy (aobolensk) <details> <summary>Changes</summary> Pointers to array elements lost the address space of the array, which broke the CIR verifier on spirv64-amd-amdhsa --- Full diff: https://github.com/llvm/llvm-project/pull/231367.diff 9 Files Affected: - (modified) clang/lib/CIR/CodeGen/CIRGenClass.cpp (+2-1) - (modified) clang/lib/CIR/CodeGen/CIRGenDecl.cpp (+4-2) - (modified) clang/lib/CIR/CodeGen/CIRGenItaniumCXXABI.cpp (+1-5) - (modified) clang/lib/CIR/Dialect/IR/CIRDialect.cpp (+4) - (modified) clang/lib/CIR/Dialect/Transforms/CXXABILowering.cpp (+5-6) - (modified) clang/lib/CIR/Dialect/Transforms/TargetLowering/CIRCXXABI.cpp (+4-2) - (modified) clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerItaniumCXXABI.cpp (+4-2) - (added) clang/test/CIR/CodeGenHIP/array-structor-addrspace.hip (+109) - (modified) clang/test/CIR/IR/invalid-array-structor.cir (+30) ``````````diff 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..33c6ca9396b09 100644 --- a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp @@ -1117,7 +1117,7 @@ struct IrregularPartialArrayDestroy final : EHScopeStack::Cleanup { QualType baseElementType = cgf.getContext().getBaseElementType(elementType); mlir::Type cirElementType = cgf.convertTypeForMem(baseElementType); - cir::PointerType ptrToElmType = builder.getPointerTo(cirElementType); + auto ptrToElmType = mlir::cast<cir::PointerType>(begin.getType()); mlir::Value begin = arrayBegin; if (baseElementType != elementType) { @@ -1270,7 +1270,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@+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@+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}> `````````` </details> https://github.com/llvm/llvm-project/pull/231367 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
