Update LLVM SHA (#541)

* Update LLVM SHA

* BugFix: Use ShapedType::kDynamic instead of magic number "-1" for "?"

Co-authored-by: [email protected] <Nishant Patel>
Co-authored-by: Sang Ik Lee <[email protected]>

Author: nbpatel

build_tools/llvm_version.txt

@@ -1 +1 @@
-06e972ed91e6d173025dc122d202f546d1a5e8ce
+ac1ec9e2904a696e360b40572c3b3c29d67981ef

lib/Conversion/DistToStandard/DistToStandard.cpp

@@ -85,15 +85,23 @@ struct RuntimePrototypesOpConverter
     auto dtypeType = rewriter.getIntegerType(sizeof(int) * 8);
     auto opType =
         rewriter.getIntegerType(sizeof(::imex::ptensor::ReduceOpId) * 8);
-    requireFunc(loc, rewriter, module, "_idtr_init_dtensor",
-                {::mlir::RankedTensorType::get({-1}, dtype), i64Type},
-                {i64Type});
-    requireFunc(loc, rewriter, module, "_idtr_local_shape",
-                {i64Type, ::mlir::RankedTensorType::get({-1}, dtype), i64Type},
-                {});
-    requireFunc(loc, rewriter, module, "_idtr_local_offsets",
-                {i64Type, ::mlir::RankedTensorType::get({-1}, dtype), i64Type},
-                {});
+    requireFunc(
+        loc, rewriter, module, "_idtr_init_dtensor",
+        {::mlir::RankedTensorType::get({::mlir::ShapedType::kDynamic}, dtype),
+         i64Type},
+        {i64Type});
+    requireFunc(
+        loc, rewriter, module, "_idtr_local_shape",
+        {i64Type,
+         ::mlir::RankedTensorType::get({::mlir::ShapedType::kDynamic}, dtype),
+         i64Type},
+        {});
+    requireFunc(
+        loc, rewriter, module, "_idtr_local_offsets",
+        {i64Type,
+         ::mlir::RankedTensorType::get({::mlir::ShapedType::kDynamic}, dtype),
+         i64Type},
+        {});
     requireFunc(loc, rewriter, module, "_idtr_reduce_all",
                 {::mlir::RankedTensorType::get({}, dtype), dtypeType, opType},
                 {});
@@ -117,10 +125,14 @@ struct RegisterPTensorOpConverter
     auto i64Rank = rewriter.create<::mlir::arith::IndexCastOp>(
         loc, rewriter.getI64Type(), rank);
     auto shapeTnsr = rewriter.create<::mlir::tensor::CastOp>(
-        loc, ::mlir::RankedTensorType::get({-1}, rewriter.getIndexType()),
+        loc,
+        ::mlir::RankedTensorType::get({::mlir::ShapedType::kDynamic},
+                                      rewriter.getIndexType()),
         shape);
     auto i64Tnsr = rewriter.create<::mlir::arith::IndexCastOp>(
-        loc, ::mlir::RankedTensorType::get({-1}, rewriter.getI64Type()),
+        loc,
+        ::mlir::RankedTensorType::get({::mlir::ShapedType::kDynamic},
+                                      rewriter.getI64Type()),
         shapeTnsr);
     auto fsa = rewriter.getStringAttr("_idtr_init_dtensor");
     rewriter.replaceOpWithNewOp<::mlir::func::CallOp>(

lib/Dialect/PTensor/Transforms/PTensorDist.cpp

@@ -189,9 +189,10 @@ struct DistARange : public RecOpRewritePattern<::imex::ptensor::ARangeOp> {
     auto stop = rewriter.create<::mlir::arith::AddIOp>(
         loc, start, tmp2); // start + (lShape[0] * stride)
     //  get type of local tensor
-    ::llvm::ArrayRef<int64_t> lShape({-1});
+    ::llvm::ArrayRef<int64_t> lShape({::mlir::ShapedType::kDynamic});
     auto artype = ::imex::ptensor::PTensorType::get(
-        rewriter.getContext(), ::mlir::RankedTensorType::get({-1}, dtype),
+        rewriter.getContext(),
+        ::mlir::RankedTensorType::get({::mlir::ShapedType::kDynamic}, dtype),
         false, false);
     // finally create local arange
     auto dmy = ::mlir::Value(); // createInt<1>(loc, rewriter, 0);

lib/Transforms/SetSPIRVCapabilities.cpp

@@ -83,9 +83,9 @@ struct SetSPIRVCapabilitiesPass
       auto triple = spirv::VerCapExtAttr::get(
           spirv::Version::V_1_0, caps_opencl, exts_opencl, context);
       auto attr = spirv::TargetEnvAttr::get(
-          triple, spirv::Vendor::Unknown, spirv::DeviceType::Unknown,
-          spirv::TargetEnvAttr::kUnknownDeviceID,
-          spirv::getDefaultResourceLimits(context));
+          triple, spirv::getDefaultResourceLimits(context),
+          spirv::ClientAPI::OpenCL, spirv::Vendor::Unknown,
+          spirv::DeviceType::Unknown, spirv::TargetEnvAttr::kUnknownDeviceID);
       auto op = getOperation();
       op->walk([&](mlir::gpu::GPUModuleOp op) {
         op->setAttr(spirv::getTargetEnvAttrName(), attr);
@@ -94,9 +94,9 @@ struct SetSPIRVCapabilitiesPass
       auto triple = spirv::VerCapExtAttr::get(
           spirv::Version::V_1_0, caps_vulkan, exts_vulkan, context);
       auto attr = spirv::TargetEnvAttr::get(
-          triple, spirv::Vendor::Unknown, spirv::DeviceType::Unknown,
-          spirv::TargetEnvAttr::kUnknownDeviceID,
-          spirv::getDefaultResourceLimits(context));
+          triple, spirv::getDefaultResourceLimits(context),
+          spirv::ClientAPI::Vulkan, spirv::Vendor::Unknown,
+          spirv::DeviceType::Unknown, spirv::TargetEnvAttr::kUnknownDeviceID);
       auto op = getOperation();
       op->walk([&](mlir::gpu::GPUModuleOp op) {
         op->setAttr(spirv::getTargetEnvAttrName(), attr);

test/Transforms/set-spirv-capability.mlir

@@ -4,9 +4,9 @@
 module attributes {gpu.container_module} {
 
 // OPENCL: module attributes {gpu.container_module} {
-// OPENCL: gpu.module @main_kernel attributes {spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Addresses, Float16Buffer, Int64, Int16, Int8, Kernel, Linkage, Vector16, GenericPointer, Groups, Float16, Float64, AtomicFloat32AddEXT, ExpectAssumeKHR], [SPV_EXT_shader_atomic_float_add, SPV_KHR_expect_assume]>, #spirv.resource_limits<>>} {
+// OPENCL: gpu.module @main_kernel attributes {spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Addresses, Float16Buffer, Int64, Int16, Int8, Kernel, Linkage, Vector16, GenericPointer, Groups, Float16, Float64, AtomicFloat32AddEXT, ExpectAssumeKHR], [SPV_EXT_shader_atomic_float_add, SPV_KHR_expect_assume]>, api=OpenCL, #spirv.resource_limits<>>} {
 // VULKAN: module attributes {gpu.container_module} {
-// VULKAN: gpu.module @main_kernel attributes {spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>>} {
+// VULKAN: gpu.module @main_kernel attributes {spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, api=Vulkan, #spirv.resource_limits<>>} {
   gpu.module @main_kernel {
     gpu.func @main_kernel(%arg0: memref<8xf32>, %arg1: memref<8xf32>, %arg2: memref<8xf32>) kernel attributes {spirv.entry_point_abi = #spirv.entry_point_abi<>} {
       cf.br ^bb1