Merged main:5434b04234a86273d719f96e9efa6332a53c0eba into amd-gfx:17b97598a1e4

Local branch amd-gfx 17b9759 Merged main:430729d71e4e582aa04b5d30796de14404fed56a into amd-gfx:cccd99d0cdce
Remote branch main 5434b04 [flang][debug] Add support for fixed size arrays. (llvm#92568)

Author: SC llvm team

flang/lib/Optimizer/Transforms/DebugTypeGenerator.cpp

@@ -37,6 +37,45 @@ static mlir::LLVM::DITypeAttr genPlaceholderType(mlir::MLIRContext *context) {
                       llvm::dwarf::DW_ATE_signed);
 }
 
+mlir::LLVM::DITypeAttr DebugTypeGenerator::convertSequenceType(
+    fir::SequenceType seqTy, mlir::LLVM::DIFileAttr fileAttr,
+    mlir::LLVM::DIScopeAttr scope, mlir::Location loc) {
+
+  mlir::MLIRContext *context = module.getContext();
+  // FIXME: Only fixed sizes arrays handled at the moment.
+  if (seqTy.hasDynamicExtents())
+    return genPlaceholderType(context);
+
+  llvm::SmallVector<mlir::LLVM::DINodeAttr> elements;
+  mlir::LLVM::DITypeAttr elemTy =
+      convertType(seqTy.getEleTy(), fileAttr, scope, loc);
+
+  for (fir::SequenceType::Extent dim : seqTy.getShape()) {
+    auto intTy = mlir::IntegerType::get(context, 64);
+    // FIXME: Only supporting lower bound of 1 at the moment. The
+    // 'SequenceType' has information about the shape but not the shift. In
+    // cases where the conversion originated during the processing of
+    // 'DeclareOp', it may be possible to pass on this information. But the
+    // type conversion should ideally be based on what information present in
+    // the type class so that it works from everywhere (e.g. when it is part
+    // of a module or a derived type.)
+    auto countAttr = mlir::IntegerAttr::get(intTy, llvm::APInt(64, dim));
+    auto lowerAttr = mlir::IntegerAttr::get(intTy, llvm::APInt(64, 1));
+    auto subrangeTy = mlir::LLVM::DISubrangeAttr::get(
+        context, countAttr, lowerAttr, nullptr, nullptr);
+    elements.push_back(subrangeTy);
+  }
+  // Apart from arrays, the `DICompositeTypeAttr` is used for other things like
+  // structure types. Many of its fields which are not applicable to arrays
+  // have been set to some valid default values.
+
+  return mlir::LLVM::DICompositeTypeAttr::get(
+      context, llvm::dwarf::DW_TAG_array_type, /*recursive id*/ {},
+      /* name */ nullptr, /* file */ nullptr, /* line */ 0, /* scope */ nullptr,
+      elemTy, mlir::LLVM::DIFlags::Zero, /* sizeInBits */ 0,
+      /*alignInBits*/ 0, elements);
+}
+
 mlir::LLVM::DITypeAttr
 DebugTypeGenerator::convertType(mlir::Type Ty, mlir::LLVM::DIFileAttr fileAttr,
                                 mlir::LLVM::DIScopeAttr scope,
@@ -69,6 +108,8 @@ DebugTypeGenerator::convertType(mlir::Type Ty, mlir::LLVM::DIFileAttr fileAttr,
     }
     return genBasicType(context, mlir::StringAttr::get(context, "complex"),
                         bitWidth * 2, llvm::dwarf::DW_ATE_complex_float);
+  } else if (auto seqTy = mlir::dyn_cast_or_null<fir::SequenceType>(Ty)) {
+    return convertSequenceType(seqTy, fileAttr, scope, loc);
   } else {
     // FIXME: These types are currently unhandled. We are generating a
     // placeholder type to allow us to test supported bits.

flang/lib/Optimizer/Transforms/DebugTypeGenerator.h

@@ -31,6 +31,10 @@ class DebugTypeGenerator {
                                      mlir::Location loc);
 
 private:
+  mlir::LLVM::DITypeAttr convertSequenceType(fir::SequenceType seqTy,
+                                             mlir::LLVM::DIFileAttr fileAttr,
+                                             mlir::LLVM::DIScopeAttr scope,
+                                             mlir::Location loc);
   mlir::ModuleOp module;
   KindMapping kindMapping;
 };

flang/test/Integration/debug-fixed-array-type-2.f90

@@ -0,0 +1,43 @@
+! RUN: %flang_fc1 -emit-llvm -debug-info-kind=standalone %s -o - | FileCheck %s
+
+program mn
+
+  integer d1(3)
+  integer d2(2, 5)
+  real d3(6, 8, 7)
+
+  i8 = fn1(d1, d2, d3)
+contains
+  function fn1(a1, b1, c1) result (res)
+    integer a1(3)
+    integer b1(2, 5)
+    real c1(6, 8, 7)
+    integer res
+    res = a1(1) + b1(1,2) + c1(3, 3, 4)
+  end function
+
+end program
+
+! CHECK-DAG: ![[INT:.*]] = !DIBasicType(name: "integer", size: 32, encoding: DW_ATE_signed)
+! CHECK-DAG: ![[REAL:.*]] = !DIBasicType(name: "real", size: 32, encoding: DW_ATE_float)
+! CHECK-DAG: ![[R1:.*]] = !DISubrange(count: 3, lowerBound: 1)
+! CHECK-DAG: ![[SUB1:.*]] = !{![[R1]]}
+! CHECK-DAG: ![[D1TY:.*]] = !DICompositeType(tag: DW_TAG_array_type, baseType: ![[INT]], elements: ![[SUB1]])
+! CHECK-DAG: !DILocalVariable(name: "d1"{{.*}}type: ![[D1TY]])
+
+! CHECK-DAG: ![[R21:.*]] = !DISubrange(count: 2, lowerBound: 1)
+! CHECK-DAG: ![[R22:.*]] = !DISubrange(count: 5, lowerBound: 1)
+! CHECK-DAG: ![[SUB2:.*]] = !{![[R21]], ![[R22]]}
+! CHECK-DAG: ![[D2TY:.*]] = !DICompositeType(tag: DW_TAG_array_type, baseType: ![[INT]], elements: ![[SUB2]])
+! CHECK-DAG: !DILocalVariable(name: "d2"{{.*}}type: ![[D2TY]])
+
+! CHECK-DAG: ![[R31:.*]] = !DISubrange(count: 6, lowerBound: 1)
+! CHECK-DAG: ![[R32:.*]] = !DISubrange(count: 8, lowerBound: 1)
+! CHECK-DAG: ![[R33:.*]] = !DISubrange(count: 7, lowerBound: 1)
+! CHECK-DAG: ![[SUB3:.*]] = !{![[R31]], ![[R32]], ![[R33]]}
+! CHECK-DAG: ![[D3TY:.*]] = !DICompositeType(tag: DW_TAG_array_type, baseType: ![[REAL]], elements: ![[SUB3]])
+! CHECK-DAG: !DILocalVariable(name: "d3"{{.*}}type: ![[D3TY]])
+
+! CHECK-DAG: !DILocalVariable(name: "a1", arg: 1{{.*}}type: ![[D1TY]])
+! CHECK-DAG: !DILocalVariable(name: "b1", arg: 2{{.*}}type: ![[D2TY]])
+! CHECK-DAG: !DILocalVariable(name: "c1", arg: 3{{.*}}type: ![[D3TY]])

flang/test/Transforms/debug-fixed-array-type.fir

@@ -0,0 +1,34 @@
+// RUN: fir-opt --add-debug-info --mlir-print-debuginfo %s | FileCheck %s
+
+module attributes {} {
+  func.func @_QQmain() attributes {fir.bindc_name = "mn"} {
+    %c7 = arith.constant 7 : index
+    %c8 = arith.constant 8 : index
+    %c6 = arith.constant 6 : index
+    %c5 = arith.constant 5 : index
+    %c2 = arith.constant 2 : index
+    %c3 = arith.constant 3 : index
+    %0 = fir.alloca !fir.array<3xi32> {bindc_name = "d1", uniq_name = "_QFEd1"}
+    %1 = fircg.ext_declare %0(%c3) {uniq_name = "_QFEd1"} : (!fir.ref<!fir.array<3xi32>>, index) -> !fir.ref<!fir.array<3xi32>> loc(#loc1)
+    %2 = fir.address_of(@_QFEd2) : !fir.ref<!fir.array<2x5xi32>>
+    %3 = fircg.ext_declare %2(%c2, %c5) {uniq_name = "_QFEd2"} : (!fir.ref<!fir.array<2x5xi32>>, index, index) -> !fir.ref<!fir.array<2x5xi32>> loc(#loc2)
+    %4 = fir.address_of(@_QFEd3) : !fir.ref<!fir.array<6x8x7xf32>>
+    %5 = fircg.ext_declare %4(%c6, %c8, %c7) {uniq_name = "_QFEd3"} : (!fir.ref<!fir.array<6x8x7xf32>>, index, index, index) -> !fir.ref<!fir.array<6x8x7xf32>> loc(#loc3)
+    return
+  } loc(#loc4)
+}
+
+#loc1 = loc("test.f90":5:1)
+#loc2 = loc("test.f90":6:11)
+#loc3 = loc("test.f90":7:11)
+#loc4 = loc("test.f90":2:8)
+
+
+// CHECK-DAG: #[[INT:.*]] = #llvm.di_basic_type<tag = DW_TAG_base_type, name = "integer", sizeInBits = 32, encoding = DW_ATE_signed>
+// CHECK-DAG: #[[REAL:.*]] = #llvm.di_basic_type<tag = DW_TAG_base_type, name = "real", sizeInBits = 32, encoding = DW_ATE_float>
+// CHECK-DAG: #[[D1TY:.*]] = #llvm.di_composite_type<tag = DW_TAG_array_type{{.*}}baseType = #[[INT]], elements = #llvm.di_subrange<count = 3 : i64, lowerBound = 1 : i64>>
+// CHECK-DAG: #[[D2TY:.*]] = #llvm.di_composite_type<tag = DW_TAG_array_type{{.*}}baseType = #[[INT]], elements = #llvm.di_subrange<count = 2 : i64, lowerBound = 1 : i64>, #llvm.di_subrange<count = 5 : i64, lowerBound = 1 : i64>>
+// CHECK-DAG: #[[D3TY:.*]] = #llvm.di_composite_type<tag = DW_TAG_array_type{{.*}}baseType = #[[REAL]], elements = #llvm.di_subrange<count = 6 : i64, lowerBound = 1 : i64>, #llvm.di_subrange<count = 8 : i64, lowerBound = 1 : i64>, #llvm.di_subrange<count = 7 : i64, lowerBound = 1 : i64>>
+// CHECK-DAG: #llvm.di_local_variable<{{.*}}name = "d1"{{.*}}type = #[[D1TY]]>
+// CHECK-DAG: #llvm.di_local_variable<{{.*}}name = "d2"{{.*}}type = #[[D2TY]]>
+// CHECK-DAG: #llvm.di_local_variable<{{.*}}name = "d3"{{.*}}type = #[[D3TY]]>

llvm/include/llvm/Config/llvm-config.h.cmake

@@ -16,7 +16,7 @@
 
 /* Indicate that this is LLVM compiled from the amd-gfx branch. */
 #define LLVM_HAVE_BRANCH_AMD_GFX
-#define LLVM_MAIN_REVISION 499704
+#define LLVM_MAIN_REVISION 499708
 
 /* Define if LLVM_ENABLE_DUMP is enabled */
 #cmakedefine LLVM_ENABLE_DUMP

llvm/lib/Target/X86/X86ISelLowering.cpp

@@ -42930,7 +42930,6 @@ bool X86TargetLowering::isGuaranteedNotToBeUndefOrPoisonForTargetNode(
     bool PoisonOnly, unsigned Depth) const {
   unsigned NumElts = DemandedElts.getBitWidth();
 
-  // TODO: Add more target shuffles.
   switch (Op.getOpcode()) {
   case X86ISD::PSHUFD:
   case X86ISD::VPERMILPI: {
@@ -42966,8 +42965,12 @@ bool X86TargetLowering::canCreateUndefOrPoisonForTargetNode(
     SDValue Op, const APInt &DemandedElts, const SelectionDAG &DAG,
     bool PoisonOnly, bool ConsiderFlags, unsigned Depth) const {
 
-  // TODO: Add more target shuffles.
   switch (Op.getOpcode()) {
+  // SSE vector shifts handle out of bounds shift amounts.
+  case X86ISD::VSHLI:
+  case X86ISD::VSRLI:
+  case X86ISD::VSRAI:
+    return false;
   case X86ISD::PSHUFD:
   case X86ISD::VPERMILPI:
   case X86ISD::UNPCKH:

llvm/test/CodeGen/X86/freeze-binary.ll

@@ -546,9 +546,8 @@ define <8 x i16> @freeze_ashr_vec(<8 x i16> %a0) nounwind {
 define <4 x i32> @freeze_ashr_vec_outofrange(<4 x i32> %a0) nounwind {
 ; X86-LABEL: freeze_ashr_vec_outofrange:
 ; X86:       # %bb.0:
-; X86-NEXT:    psrad $1, %xmm0
 ; X86-NEXT:    pshufd {{.*#+}} xmm0 = xmm0[0,3,2,3]
-; X86-NEXT:    psrad $2, %xmm0
+; X86-NEXT:    psrad $3, %xmm0
 ; X86-NEXT:    retl
 ;
 ; X64-LABEL: freeze_ashr_vec_outofrange:
@@ -660,9 +659,8 @@ define <8 x i16> @freeze_lshr_vec(<8 x i16> %a0) nounwind {
 define <4 x i32> @freeze_lshr_vec_outofrange(<4 x i32> %a0) nounwind {
 ; X86-LABEL: freeze_lshr_vec_outofrange:
 ; X86:       # %bb.0:
-; X86-NEXT:    psrld $1, %xmm0
 ; X86-NEXT:    pshufd {{.*#+}} xmm0 = xmm0[0,3,2,3]
-; X86-NEXT:    psrld $2, %xmm0
+; X86-NEXT:    psrld $3, %xmm0
 ; X86-NEXT:    retl
 ;
 ; X64-LABEL: freeze_lshr_vec_outofrange:

mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.h

@@ -13,7 +13,6 @@
 #include "mlir/Dialect/LLVMIR/Transforms/LegalizeForExport.h"
 #include "mlir/Dialect/LLVMIR/Transforms/OptimizeForNVVM.h"
 #include "mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h"
-#include "mlir/Dialect/LLVMIR/Transforms/TypeConsistency.h"
 #include "mlir/Pass/Pass.h"
 
 namespace mlir {

mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td

@@ -43,24 +43,6 @@ def LLVMRequestCWrappers
   let constructor = "::mlir::LLVM::createRequestCWrappersPass()";
 }
 
-def LLVMTypeConsistency
-    : Pass<"llvm-type-consistency", "::mlir::LLVM::LLVMFuncOp"> {
-  let summary = "Rewrites to improve type consistency";
-  let description = [{
-    Set of rewrites to improve the coherency of types within an LLVM dialect
-    program. This will adjust operations operating on pointers so they interpret
-    their associated pointee type as consistently as possible.
-  }];
-  let constructor = "::mlir::LLVM::createTypeConsistencyPass()";
-
-  let options = [
-    Option<"maxVectorSplitSize", "max-vector-split-size", "unsigned",
-           /*default=*/"512",
-           "Maximum size in bits of a vector value in a load or store operation"
-           " operating on multiple elements that should still be split">,
-  ];
-}
-
 def NVVMOptimizeForTarget : Pass<"llvm-optimize-for-nvvm-target"> {
   let summary = "Optimize NVVM IR";
   let constructor = "::mlir::NVVM::createOptimizeForTargetPass()";