https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/231404
>From 546645f9883ec0264d0ae267022738cf5370b9c1 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Tue, 13 Oct 2026 10:03:23 +0200 Subject: [PATCH 1/2] [CIR] Support building vectors from smaller vectors Fix a crash on OpenCL `(float4)(a, b)` and HIP `__half2 data{xx}` Use shuffles like classic codegen --- clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp | 49 ++++++++++++++++++++++ clang/test/CIR/CodeGen/vector-ext.cpp | 28 +++++++++++++ clang/test/CIR/CodeGenOpenCL/vector.cl | 39 +++++++++++++++++ 3 files changed, 116 insertions(+) diff --git a/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp b/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp index fbf3e90d14369..3fa8f985e9c9b 100644 --- a/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp @@ -26,6 +26,7 @@ #include "mlir/IR/Value.h" #include <cassert> +#include <numeric> #include <utility> using namespace clang; @@ -420,6 +421,9 @@ class ScalarExprEmitter : public StmtVisitor<ScalarExprEmitter, mlir::Value> { } mlir::Value VisitInitListExpr(InitListExpr *e); + mlir::Value emitVectorInitWithSubvectors(mlir::Location loc, + cir::VectorType vectorType, + ArrayRef<mlir::Value> inits); mlir::Value VisitArrayInitIndexExpr(ArrayInitIndexExpr *e) { assert(cgf.getArrayInitIndex() && @@ -2948,6 +2952,45 @@ mlir::Value ScalarExprEmitter::VisitMemberExpr(MemberExpr *e) { return emitLoadOfLValue(e); } +mlir::Value +ScalarExprEmitter::emitVectorInitWithSubvectors(mlir::Location loc, + cir::VectorType vectorType, + ArrayRef<mlir::Value> inits) { + const unsigned resElts = vectorType.getSize(); + mlir::Value result = + builder.getConstant(loc, cir::PoisonAttr::get(vectorType)); + unsigned curIdx = 0; + for (mlir::Value init : inits) { + auto initType = mlir::dyn_cast<cir::VectorType>(init.getType()); + if (!initType) { + result = builder.createInsertElement(loc, result, init, curIdx++); + continue; + } + + const unsigned initElts = initType.getSize(); + SmallVector<int64_t, 16> mask(resElts, -1); + std::iota(mask.begin(), mask.begin() + initElts, 0); + mlir::Value ext = builder.createVecShuffle(loc, init, mask); + + // Poison stays on the RHS of the first shuffle, matching classic codegen. + const unsigned offset = curIdx == 0 ? 0 : resElts; + std::iota(mask.begin(), mask.begin() + curIdx, 0); + std::iota(mask.begin() + curIdx, mask.begin() + curIdx + initElts, offset); + if (curIdx == 0) + std::swap(result, ext); + result = builder.createVecShuffle(loc, result, ext, mask); + curIdx += initElts; + } + + if (curIdx < resElts) { + const mlir::Value zeroValue = + builder.getNullValue(vectorType.getElementType(), loc); + for (; curIdx < resElts; ++curIdx) + result = builder.createInsertElement(loc, result, zeroValue, curIdx); + } + return result; +} + mlir::Value ScalarExprEmitter::VisitInitListExpr(InitListExpr *e) { const unsigned numInitElements = e->getNumInits(); @@ -2970,6 +3013,12 @@ mlir::Value ScalarExprEmitter::VisitInitListExpr(InitListExpr *e) { elements.push_back(Visit(init)); } + if (llvm::any_of(elements, [](mlir::Value v) { + return mlir::isa<cir::VectorType>(v.getType()); + })) + return emitVectorInitWithSubvectors(cgf.getLoc(e->getSourceRange()), + vectorType, elements); + // Zero-initialize any remaining values. if (numInitElements < vectorType.getSize()) { const mlir::Value zeroValue = cgf.getBuilder().getNullValue( diff --git a/clang/test/CIR/CodeGen/vector-ext.cpp b/clang/test/CIR/CodeGen/vector-ext.cpp index ddde0a94eef23..e352aead72b5f 100644 --- a/clang/test/CIR/CodeGen/vector-ext.cpp +++ b/clang/test/CIR/CodeGen/vector-ext.cpp @@ -1463,3 +1463,31 @@ vi4 ternary_expression_with_vec_cond() { // OGCG: %[[COND:.*]] = icmp slt <4 x i32> %[[TMP_C]], zeroinitializer // OGCG: %[[RESULT:.*]] = select <4 x i1> %[[COND]], <4 x i32> %[[TMP_A]], <4 x i32> %[[TMP_B]] // OGCG: ret <4 x i32> %[[RESULT]] + +void init_list_same_vector(vi4 v) { vi4 w{v}; } + +// CIR-LABEL: cir.func {{.*}} @_Z21init_list_same_vectorDv4_i +// CIR: %[[V_ADDR:.*]] = cir.alloca "v" +// CIR: %[[W_ADDR:.*]] = cir.alloca "w" +// CIR: %[[V:.*]] = cir.load {{.*}} %[[V_ADDR]] : !cir.ptr<!cir.vector<4 x !s32i>>, !cir.vector<4 x !s32i> +// CIR: %[[POISON:.*]] = cir.const #cir.poison : !cir.vector<4 x !s32i> +// CIR: %[[POISON_2:.*]] = cir.const #cir.poison : !cir.vector<4 x !s32i> +// CIR: %[[EXT:.*]] = cir.vec.shuffle(%[[V]], %[[POISON_2]] : !cir.vector<4 x !s32i>) [#cir.int<0> : !s32i, #cir.int<1> : !s32i, #cir.int<2> : !s32i, #cir.int<3> : !s32i] : !cir.vector<4 x !s32i> +// CIR: %[[RESULT:.*]] = cir.vec.shuffle(%[[EXT]], %[[POISON]] : !cir.vector<4 x !s32i>) [#cir.int<0> : !s32i, #cir.int<1> : !s32i, #cir.int<2> : !s32i, #cir.int<3> : !s32i] : !cir.vector<4 x !s32i> +// CIR: cir.store {{.*}} %[[RESULT]], %[[W_ADDR]] : !cir.vector<4 x !s32i>, !cir.ptr<!cir.vector<4 x !s32i>> + +// LLVM-LABEL: define {{.*}} @_Z21init_list_same_vectorDv4_i +// LLVM: %[[V_ADDR:.*]] = alloca <4 x i32> +// LLVM: %[[W_ADDR:.*]] = alloca <4 x i32> +// LLVM: %[[V:.*]] = load <4 x i32>, ptr %[[V_ADDR]], align 16 +// LLVM: %[[EXT:.*]] = shufflevector <4 x i32> %[[V]], <4 x i32> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 3> +// LLVM: %[[RESULT:.*]] = shufflevector <4 x i32> %[[EXT]], <4 x i32> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 3> +// LLVM: store <4 x i32> %[[RESULT]], ptr %[[W_ADDR]], align 16 + +// OGCG-LABEL: define {{.*}} @_Z21init_list_same_vectorDv4_i +// OGCG: %[[V_ADDR:.*]] = alloca <4 x i32> +// OGCG: %[[W_ADDR:.*]] = alloca <4 x i32> +// OGCG: %[[V:.*]] = load <4 x i32>, ptr %[[V_ADDR]], align 16 +// OGCG: %[[EXT:.*]] = shufflevector <4 x i32> %[[V]], <4 x i32> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 3> +// OGCG: %[[RESULT:.*]] = shufflevector <4 x i32> %[[EXT]], <4 x i32> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 3> +// OGCG: store <4 x i32> %[[RESULT]], ptr %[[W_ADDR]], align 16 diff --git a/clang/test/CIR/CodeGenOpenCL/vector.cl b/clang/test/CIR/CodeGenOpenCL/vector.cl index 4013449cb49b2..dafb700821a77 100644 --- a/clang/test/CIR/CodeGenOpenCL/vector.cl +++ b/clang/test/CIR/CodeGenOpenCL/vector.cl @@ -7,6 +7,7 @@ // RUN: %clang_cc1 %s -emit-llvm -triple spir-unknown-unknown -o %t.ll // RUN: FileCheck %s --input-file=%t.ll --check-prefix=LLVM +typedef __attribute__(( ext_vector_type(2) )) int int2; typedef __attribute__(( ext_vector_type(4) )) int int4; typedef __attribute__(( ext_vector_type(4) )) float float4; @@ -104,3 +105,41 @@ int4 vec_unary_dec(int4 a) { // LLVM: %[[VEC_DEC:.*]] = add <4 x i32> %[[A:.*]], splat (i32 -1) // LLVM: ret <4 x i32> %[[VEC_DEC]] + +int4 vec_init_subvectors(int2 a, int2 b) { + return (int4)(a, b); +} + +// CIR-LABEL: cir.func {{.*}} @vec_init_subvectors +// CIR: %[[A:.*]] = cir.load {{.*}} : !cir.ptr<!cir.vector<2 x !s32i>>, !cir.vector<2 x !s32i> +// CIR: %[[B:.*]] = cir.load {{.*}} : !cir.ptr<!cir.vector<2 x !s32i>>, !cir.vector<2 x !s32i> +// CIR: %[[POISON:.*]] = cir.const #cir.poison : !cir.vector<4 x !s32i> +// CIR: %[[A_EXT:.*]] = cir.vec.shuffle(%[[A]], %{{.*}} : !cir.vector<2 x !s32i>) [#cir.int<0> : !s32i, #cir.int<1> : !s32i, #cir.int<-1> : !s32i, #cir.int<-1> : !s32i] : !cir.vector<4 x !s32i> +// CIR: %[[V0:.*]] = cir.vec.shuffle(%[[A_EXT]], %[[POISON]] : !cir.vector<4 x !s32i>) [#cir.int<0> : !s32i, #cir.int<1> : !s32i, #cir.int<-1> : !s32i, #cir.int<-1> : !s32i] : !cir.vector<4 x !s32i> +// CIR: %[[B_EXT:.*]] = cir.vec.shuffle(%[[B]], %{{.*}} : !cir.vector<2 x !s32i>) [#cir.int<0> : !s32i, #cir.int<1> : !s32i, #cir.int<-1> : !s32i, #cir.int<-1> : !s32i] : !cir.vector<4 x !s32i> +// CIR: cir.vec.shuffle(%[[V0]], %[[B_EXT]] : !cir.vector<4 x !s32i>) [#cir.int<0> : !s32i, #cir.int<1> : !s32i, #cir.int<4> : !s32i, #cir.int<5> : !s32i] : !cir.vector<4 x !s32i> + +// LLVM-LABEL: @vec_init_subvectors +// LLVM: %[[RESULT:.*]] = shufflevector <2 x i32> %{{.*}}, <2 x i32> %{{.*}}, <4 x i32> <i32 0, i32 1, i32 2, i32 3> +// LLVM: ret <4 x i32> %[[RESULT]] + +int4 vec_init_scalar_and_subvector(int2 a, int c) { + return (int4)(c, a, c); +} + +// CIR-LABEL: cir.func {{.*}} @vec_init_scalar_and_subvector +// CIR: %[[C:.*]] = cir.load {{.*}} : !cir.ptr<!s32i>, !s32i +// CIR: %[[A:.*]] = cir.load {{.*}} : !cir.ptr<!cir.vector<2 x !s32i>>, !cir.vector<2 x !s32i> +// CIR: %[[C2:.*]] = cir.load {{.*}} : !cir.ptr<!s32i>, !s32i +// CIR: %[[POISON:.*]] = cir.const #cir.poison : !cir.vector<4 x !s32i> +// CIR: %[[V0:.*]] = cir.vec.insert %[[C]], %[[POISON]][%{{.*}} : !u64i] : !cir.vector<4 x !s32i> +// CIR: %[[A_EXT:.*]] = cir.vec.shuffle(%[[A]], %{{.*}} : !cir.vector<2 x !s32i>) [#cir.int<0> : !s32i, #cir.int<1> : !s32i, #cir.int<-1> : !s32i, #cir.int<-1> : !s32i] : !cir.vector<4 x !s32i> +// CIR: %[[V1:.*]] = cir.vec.shuffle(%[[V0]], %[[A_EXT]] : !cir.vector<4 x !s32i>) [#cir.int<0> : !s32i, #cir.int<4> : !s32i, #cir.int<5> : !s32i, #cir.int<-1> : !s32i] : !cir.vector<4 x !s32i> +// CIR: cir.vec.insert %[[C2]], %[[V1]][%{{.*}} : !u64i] : !cir.vector<4 x !s32i> + +// LLVM-LABEL: @vec_init_scalar_and_subvector +// LLVM: %[[V0:.*]] = insertelement <4 x i32> poison, i32 %[[C:.*]], i64 0 +// LLVM: %[[EXT:.*]] = shufflevector <2 x i32> %{{.*}}, <2 x i32> poison, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison> +// LLVM: %[[V1:.*]] = shufflevector <4 x i32> %[[V0]], <4 x i32> %[[EXT]], <4 x i32> <i32 0, i32 4, i32 5, i32 poison> +// LLVM: %[[V2:.*]] = insertelement <4 x i32> %[[V1]], i32 %[[C]], i64 3 +// LLVM: ret <4 x i32> %[[V2]] >From f84ce4b67ecddba7b00be4ae44eb1ff6a87be157 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Wed, 14 Oct 2026 11:02:58 +0200 Subject: [PATCH 2/2] Address comments --- clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp b/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp index 3fa8f985e9c9b..0a2c171bc2161 100644 --- a/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp @@ -2968,6 +2968,8 @@ ScalarExprEmitter::emitVectorInitWithSubvectors(mlir::Location loc, } const unsigned initElts = initType.getSize(); + assert(curIdx + initElts <= resElts && + "Sema should reject excess elements"); SmallVector<int64_t, 16> mask(resElts, -1); std::iota(mask.begin(), mask.begin() + initElts, 0); mlir::Value ext = builder.createVecShuffle(loc, init, mask); @@ -3009,13 +3011,13 @@ mlir::Value ScalarExprEmitter::VisitInitListExpr(InitListExpr *e) { mlir::cast<cir::VectorType>(cgf.convertType(e->getType())); SmallVector<mlir::Value, 16> elements; + bool hasVectorInit = false; for (Expr *init : e->inits()) { elements.push_back(Visit(init)); + hasVectorInit |= mlir::isa<cir::VectorType>(elements.back().getType()); } - if (llvm::any_of(elements, [](mlir::Value v) { - return mlir::isa<cir::VectorType>(v.getType()); - })) + if (hasVectorInit) return emitVectorInitWithSubvectors(cgf.getLoc(e->getSourceRange()), vectorType, elements); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
