llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clangir
@llvm/pr-subscribers-clang
Author: Arseniy Obolenskiy (aobolensk)
<details>
<summary>Changes</summary>
Fix a crash on OpenCL `(float4)(a, b)` and HIP `__half2 data{xx}`
Use shuffles like classic codegen
---
Full diff: https://github.com/llvm/llvm-project/pull/231404.diff
3 Files Affected:
- (modified) clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp (+49)
- (modified) clang/test/CIR/CodeGen/vector-ext.cpp (+28)
- (modified) clang/test/CIR/CodeGenOpenCL/vector.cl (+39)
``````````diff
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]]
``````````
</details>
https://github.com/llvm/llvm-project/pull/231404
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
