Author: Arseniy Obolenskiy Date: 2026-10-30T16:51:35+01:00 New Revision: 4c5f3a3179a10c4d85c694d40aec336ec6098165
URL: https://github.com/llvm/llvm-project/commit/4c5f3a3179a10c4d85c694d40aec336ec6098165 DIFF: https://github.com/llvm/llvm-project/commit/4c5f3a3179a10c4d85c694d40aec336ec6098165.diff LOG: [CIR] Keep the address space in dynamic_cast to void (#237245) Use the source address space instead of the default one This fixes CIR verification for HIP on spirv64-amd-amdhsa Added: clang/test/CIR/CodeGenHIP/dynamic-cast-void-addrspace.hip Modified: clang/lib/CIR/CodeGen/CIRGenBuilder.h clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerItaniumCXXABI.cpp Removed: ################################################################################ diff --git a/clang/lib/CIR/CodeGen/CIRGenBuilder.h b/clang/lib/CIR/CodeGen/CIRGenBuilder.h index e60e3092b699f3..48e657aa2bfd6d 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuilder.h +++ b/clang/lib/CIR/CodeGen/CIRGenBuilder.h @@ -468,9 +468,8 @@ class CIRGenBuilderTy : public cir::CIRBaseBuilderTy { mlir::Value createDynCastToVoid(mlir::Location loc, mlir::Value src, bool vtableUseRelativeLayout) { - // TODO(cir): consider address space here. - assert(!cir::MissingFeatures::addressSpace()); - cir::PointerType destTy = getVoidPtrTy(); + auto srcPtrTy = mlir::cast<cir::PointerType>(src.getType()); + cir::PointerType destTy = getVoidPtrTy(srcPtrTy.getAddrSpace()); return cir::DynamicCastOp::create( *this, loc, destTy, cir::DynamicCastKind::Ptr, src, cir::DynamicCastInfoAttr{}, vtableUseRelativeLayout); diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerItaniumCXXABI.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerItaniumCXXABI.cpp index 962cdd3782adad..f5179f38cfc1e4 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerItaniumCXXABI.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerItaniumCXXABI.cpp @@ -789,7 +789,7 @@ static mlir::Value buildDynamicCastToVoidAfterNullCheck( mlir::Location loc = op.getLoc(); bool vtableUsesRelativeLayout = op.getRelativeLayout(); - // TODO(cir): consider address space in this function. + // TODO(cir): consider address space of the vtable pointer. assert(!cir::MissingFeatures::addressSpace()); mlir::Type vtableElemTy; @@ -837,20 +837,18 @@ static mlir::Value buildDynamicCastToVoidAfterNullCheck( /*mem_order=*/cir::MemOrderAttr(), /*invariant=*/false); - auto voidPtrTy = - cir::PointerType::get(cir::VoidType::get(builder.getContext())); - // Add the offset to the given pointer to get the cast result. // Cast the input pointer to a uint8_t* to allow pointer arithmetic. mlir::Type u8PtrTy = cir::PointerType::get(cir::IntType::get(builder.getContext(), /*width=*/8, - /*isSigned=*/false)); + /*isSigned=*/false), + op.getSrc().getType().getAddrSpace()); mlir::Value srcBytePtr = cir::CastOp::create( builder, loc, u8PtrTy, cir::CastKind::bitcast, op.getSrc()); auto dstBytePtr = cir::PtrStrideOp::create(builder, loc, u8PtrTy, srcBytePtr, offsetToTop); // Cast the result to a void*. - return cir::CastOp::create(builder, loc, voidPtrTy, cir::CastKind::bitcast, + return cir::CastOp::create(builder, loc, op.getType(), cir::CastKind::bitcast, dstBytePtr); } diff --git a/clang/test/CIR/CodeGenHIP/dynamic-cast-void-addrspace.hip b/clang/test/CIR/CodeGenHIP/dynamic-cast-void-addrspace.hip new file mode 100644 index 00000000000000..c52719fbd4ab0d --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/dynamic-cast-void-addrspace.hip @@ -0,0 +1,14 @@ +// 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 + +#define __device__ __attribute__((device)) + +struct A { __device__ virtual void f(); }; + +// CHECK-LABEL: cir.func {{.*}} @_Z9cast_voidP1A +// CHECK: %[[SRC:.*]] = cir.load {{.*}}, !cir.ptr<!rec_A, target_address_space(4)> +// CHECK: %[[RAW:.*]] = cir.cast bitcast %[[SRC]] : !cir.ptr<!rec_A, target_address_space(4)> -> !cir.ptr<!u8i, target_address_space(4)> +// CHECK: %[[ADJ:.*]] = cir.ptr_stride %[[RAW]], %{{.*}} : (!cir.ptr<!u8i, target_address_space(4)>, !s64i) -> !cir.ptr<!u8i, target_address_space(4)> +// CHECK: cir.cast bitcast %[[ADJ]] : !cir.ptr<!u8i, target_address_space(4)> -> !cir.ptr<!void, target_address_space(4)> +// CHECK: }) : (!cir.bool) -> !cir.ptr<!void, target_address_space(4)> +__device__ void *cast_void(A *p) { return dynamic_cast<void *>(p); } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
