[NFC] Simplify repeated vector push backs (#16936)
Simplify some verbose sequences of vector push backs. Depending on the
context:
* I left some sequences of push backs as-is, as having one per line made
it easier to make sure that the appended values matches some spec (e.g.,
ABI, linker flags).
* Turned some into direct vector initialization.
* Turned appends of repeated values into `.append(n, value)`
* Turned repeated push backs into `append_values`.
diff --git a/compiler/plugins/input/StableHLO/Conversion/StableHLOToLinalg.cpp b/compiler/plugins/input/StableHLO/Conversion/StableHLOToLinalg.cpp
index da4d752..32b54da 100644
--- a/compiler/plugins/input/StableHLO/Conversion/StableHLOToLinalg.cpp
+++ b/compiler/plugins/input/StableHLO/Conversion/StableHLOToLinalg.cpp
@@ -2156,11 +2156,9 @@
auto scatterInputMap =
AffineMap::get(broadcastShape.size(), /*symbolCount=*/0, broadcastExprs,
b.getContext());
- SmallVector<AffineMap> scatterIndexingMaps;
- scatterIndexingMaps.push_back(scatterInputMap);
- scatterIndexingMaps.push_back(scatterInputMap);
- scatterIndexingMaps.push_back(
- b.getMultiDimIdentityMap(broadcastShape.size()));
+ SmallVector<AffineMap> scatterIndexingMaps = {
+ scatterInputMap, scatterInputMap,
+ b.getMultiDimIdentityMap(broadcastShape.size())};
auto scatterGeneric = b.create<linalg::GenericOp>(
/*resultTensors=*/ArrayRef<Type>{scatterFill.getType()},
diff --git a/compiler/src/iree/compiler/Codegen/Common/ConcretizePadResultShape.cpp b/compiler/src/iree/compiler/Codegen/Common/ConcretizePadResultShape.cpp
index ca7b79b..cde65cd 100644
--- a/compiler/src/iree/compiler/Codegen/Common/ConcretizePadResultShape.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/ConcretizePadResultShape.cpp
@@ -73,12 +73,11 @@
bindSymbols(context, sym0, sym1, sym2);
auto addMap = AffineMap::get(0, 3, {sym0 + sym1 + sym2}, context);
- SmallVector<Value, 3> valueSizes;
for (int dimIndex = 0; dimIndex < rank; ++dimIndex) {
- valueSizes.clear();
- valueSizes.push_back(getAsIndexValue(lowPad[dimIndex], rewriter, loc));
- valueSizes.push_back(getAsIndexValue(source[dimIndex], rewriter, loc));
- valueSizes.push_back(getAsIndexValue(highPad[dimIndex], rewriter, loc));
+ SmallVector<Value, 3> valueSizes = {
+ getAsIndexValue(lowPad[dimIndex], rewriter, loc),
+ getAsIndexValue(source[dimIndex], rewriter, loc),
+ getAsIndexValue(highPad[dimIndex], rewriter, loc)};
// The pad op's result shape is low padding + source size + high padding.
// Try to see if we can get a constant number by composing and
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp b/compiler/src/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp
index 4cc066a..7bf0d1a 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp
@@ -148,10 +148,11 @@
tileSizes.push_back(
TileWorkgroupSizePair({{128, 256, 16}, {128, 2, 1}, 4}));
}
- tileSizes.push_back(TileWorkgroupSizePair({{32, 32, 16}, {64, 2, 1}, 4}));
- tileSizes.push_back(TileWorkgroupSizePair({{16, 32, 16}, {64, 1, 1}, 4}));
- tileSizes.push_back(TileWorkgroupSizePair({{32, 16, 16}, {32, 2, 1}, 4}));
- tileSizes.push_back(TileWorkgroupSizePair({{16, 16, 16}, {32, 1, 1}, 4}));
+ llvm::append_values(tileSizes,
+ TileWorkgroupSizePair({{32, 32, 16}, {64, 2, 1}, 4}),
+ TileWorkgroupSizePair({{16, 32, 16}, {64, 1, 1}, 4}),
+ TileWorkgroupSizePair({{32, 16, 16}, {32, 2, 1}, 4}),
+ TileWorkgroupSizePair({{16, 16, 16}, {32, 1, 1}, 4}));
}
}
diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/KernelConfig.cpp b/compiler/src/iree/compiler/Codegen/SPIRV/KernelConfig.cpp
index 7bbd45f..f053c0f 100644
--- a/compiler/src/iree/compiler/Codegen/SPIRV/KernelConfig.cpp
+++ b/compiler/src/iree/compiler/Codegen/SPIRV/KernelConfig.cpp
@@ -770,9 +770,8 @@
workgroupTileSizes.resize(lastParallelDim + 1);
threadTileSizes.resize(lastParallelDim + 1);
- tileSizes.push_back(workgroupTileSizes);
- tileSizes.push_back(threadTileSizes);
- tileSizes.push_back(reductionTileSizes);
+ llvm::append_values(tileSizes, workgroupTileSizes, threadTileSizes,
+ reductionTileSizes);
return setOpConfigAndEntryPointFnTranslation(
op->getParentOfType<mlir::FunctionOpInterface>(), op, tileSizes,
CodeGenPipeline::SPIRVBaseVectorize, workgroupSize);
@@ -954,12 +953,8 @@
reductionTileSizes.append(kIndex, 0);
reductionTileSizes.push_back(schedule->kTileCount * schedule->kSize);
- TileSizesListType tileSizes;
- tileSizes.reserve(3);
- tileSizes.push_back(workgroupTileSizes);
- tileSizes.push_back(subgroupTileSizes);
- tileSizes.push_back(reductionTileSizes);
- tileSizes.push_back(vectorSizes);
+ TileSizesListType tileSizes = {workgroupTileSizes, subgroupTileSizes,
+ reductionTileSizes, vectorSizes};
// Don't do multibuffering if the inner reduction loop is folded out.
auto pipelineDepth = softwarePipelineDepth;
diff --git a/compiler/src/iree/compiler/Codegen/TransformStrategies/Common/Common.cpp b/compiler/src/iree/compiler/Codegen/TransformStrategies/Common/Common.cpp
index 50eff87..c12ff9a 100644
--- a/compiler/src/iree/compiler/Codegen/TransformStrategies/Common/Common.cpp
+++ b/compiler/src/iree/compiler/Codegen/TransformStrategies/Common/Common.cpp
@@ -279,10 +279,10 @@
SmallVector<Attribute> transposeAttrs;
for (auto &transp : transposePaddings)
transposeAttrs.push_back(b.getI64ArrayAttr(transp));
- SmallVector<Type> resultTypes;
- resultTypes.push_back(opH.getType());
- resultTypes.push_back(transform::AnyOpType::get(b.getContext()));
- resultTypes.push_back(transform::AnyOpType::get(b.getContext()));
+
+ Type resultTypes[] = {opH.getType(),
+ transform::AnyOpType::get(b.getContext()),
+ transform::AnyOpType::get(b.getContext())};
return b
.create<transform::PadOp>(resultTypes, opH, b.getArrayAttr(paddingValues),
b.getI64ArrayAttr(paddingDimensions),
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertCommandBufferOps.cpp b/compiler/src/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertCommandBufferOps.cpp
index cd7fdc8..4de1867 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertCommandBufferOps.cpp
+++ b/compiler/src/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertCommandBufferOps.cpp
@@ -8,6 +8,7 @@
#include "iree/compiler/Dialect/Util/IR/UtilTypes.h"
#include "iree/compiler/Dialect/VM/Conversion/ImportUtils.h"
#include "iree/compiler/Dialect/VM/IR/VMOps.h"
+#include "llvm/ADT/STLExtras.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Transforms/DialectConversion.h"
@@ -171,7 +172,7 @@
// %recv_offset : i64,
// %recv_length : i64,
// %element_count : i64
- SmallVector<Value, 8> callOperands;
+ SmallVector<Value> callOperands;
callOperands.push_back(adaptor.getCommandBuffer());
callOperands.push_back(adaptor.getChannel());
callOperands.push_back(rewriter.create<IREE::VM::ConstI32Op>(
diff --git a/compiler/src/iree/compiler/Dialect/Util/Transforms/PropagateSubranges.cpp b/compiler/src/iree/compiler/Dialect/Util/Transforms/PropagateSubranges.cpp
index f8d1cbf..3c219bd 100644
--- a/compiler/src/iree/compiler/Dialect/Util/Transforms/PropagateSubranges.cpp
+++ b/compiler/src/iree/compiler/Dialect/Util/Transforms/PropagateSubranges.cpp
@@ -13,6 +13,7 @@
#include "iree/compiler/Dialect/Util/Transforms/Patterns.h"
#include "iree/compiler/Utils/IndexSet.h"
#include "llvm/ADT/BreadthFirstIterator.h"
+#include "llvm/ADT/STLExtras.h"
#include "llvm/ADT/TypeSwitch.h"
#include "llvm/Support/Debug.h"
#include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h"
@@ -114,10 +115,8 @@
static void expandType(Type type, SmallVectorImpl<Type> &newTypes) {
newTypes.push_back(type);
if (isResourceType(type)) {
- auto indexType = IndexType::get(type.getContext());
- newTypes.push_back(indexType); // resource size
- newTypes.push_back(indexType); // subrange offset
- newTypes.push_back(indexType); // subrange length
+ // resource size, subrance offset, subrange length
+ newTypes.append(3, IndexType::get(type.getContext()));
}
}
@@ -188,10 +187,8 @@
if (isResourceType(operand.getType())) {
auto subrange =
consumeSubrange(loc, operand, subrangeMap, indexSet, builder);
- newOperands.push_back(subrange.resource);
- newOperands.push_back(subrange.resourceSize);
- newOperands.push_back(subrange.subrangeOffset);
- newOperands.push_back(subrange.subrangeLength);
+ llvm::append_values(newOperands, subrange.resource, subrange.resourceSize,
+ subrange.subrangeOffset, subrange.subrangeLength);
} else {
newOperands.push_back(operand);
}
diff --git a/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/ConvertVMToEmitC.cpp b/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/ConvertVMToEmitC.cpp
index b19adf9..926402a 100644
--- a/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/ConvertVMToEmitC.cpp
+++ b/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/ConvertVMToEmitC.cpp
@@ -21,6 +21,7 @@
#include "mlir/IR/BuiltinDialect.h"
#include "mlir/IR/BuiltinOps.h"
#include "mlir/IR/IRMapping.h"
+#include "mlir/IR/MLIRContext.h"
#include "mlir/IR/Matchers.h"
#include "mlir/Pass/Pass.h"
#include "mlir/Transforms/DialectConversion.h"
@@ -3964,7 +3965,6 @@
getOperands(IREE::VM::ListAllocOp op, Adaptor adaptor,
ConversionPatternRewriter &rewriter, Type elementType,
Value containerPtr, Value allocator) const {
- SmallVector<Value> result;
auto moduleOp = op.getOperation()->getParentOfType<IREE::VM::ModuleOp>();
auto parentFuncOp =
@@ -3980,12 +3980,8 @@
}
Value capacity = adaptor.getOperands()[0];
-
- result.push_back(elementTypePtr.value());
- result.push_back(capacity);
- result.push_back(allocator);
- result.push_back(containerPtr);
-
+ SmallVector<Value> result = {elementTypePtr.value(), capacity, allocator,
+ containerPtr};
return result;
}
@@ -3995,9 +3991,6 @@
Value containerPtr, Value allocator) const {
auto ctx = op.getContext();
auto loc = op.getLoc();
-
- SmallVector<Value> result;
-
Value access =
rewriter
.create<emitc::ConstantOp>(
@@ -4012,12 +4005,8 @@
Value length = adaptor.getOperands()[0];
Value alignment = adaptor.getOperands()[1];
- result.push_back(access);
- result.push_back(length);
- result.push_back(alignment);
- result.push_back(allocator);
- result.push_back(containerPtr);
-
+ SmallVector<Value> result = {access, length, alignment, allocator,
+ containerPtr};
return result;
}
@@ -4028,8 +4017,6 @@
auto ctx = op.getContext();
auto loc = op.getLoc();
- SmallVector<Value> result;
-
Value access =
rewriter
.create<emitc::ConstantOp>(
@@ -4061,14 +4048,8 @@
Value length = adaptor.getOperands()[2];
Value alignment = adaptor.getOperands()[3];
- result.push_back(access);
- result.push_back(source);
- result.push_back(offset);
- result.push_back(length);
- result.push_back(alignment);
- result.push_back(allocator);
- result.push_back(containerPtr);
-
+ SmallVector<Value> result = {access, source, offset, length,
+ alignment, allocator, containerPtr};
return result;
}
};
diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetEnvironment.cpp b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetEnvironment.cpp
index 541f9cc..bcf3b55 100644
--- a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetEnvironment.cpp
+++ b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetEnvironment.cpp
@@ -8,6 +8,7 @@
#include "iree/compiler/Dialect/Vulkan/IR/VulkanTypes.h"
#include "iree/compiler/Dialect/Vulkan/Utils/TargetTriple.h"
+#include "llvm/ADT/STLExtras.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVEnums.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVTypes.h"
#include "mlir/IR/Builders.h"
@@ -148,9 +149,9 @@
capabilities.push_back(spirv::Capability::VariablePointersStorageBuffer);
}
if (vkCapabilities.getShaderIntegerDotProduct()) {
- capabilities.push_back(spirv::Capability::DotProduct);
- capabilities.push_back(spirv::Capability::DotProductInputAll);
- capabilities.push_back(spirv::Capability::DotProductInput4x8BitPacked);
+ llvm::append_values(capabilities, spirv::Capability::DotProduct,
+ spirv::Capability::DotProductInputAll,
+ spirv::Capability::DotProductInput4x8BitPacked);
if (vkCapabilities.getShaderInt8()) {
capabilities.push_back(spirv::Capability::DotProductInput4x8Bit);
}
diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetTriple.cpp b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetTriple.cpp
index 54c5e73..9564bf7 100644
--- a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetTriple.cpp
+++ b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/TargetTriple.cpp
@@ -118,40 +118,31 @@
switch (triple.getArch()) {
case TargetTripleArch::Apple_M1: {
// Example: https://vulkan.gpuinfo.org/displayreport.php?id=14673
- const Extension list[] = {
- Extension::VK_KHR_16bit_storage,
- Extension::VK_KHR_8bit_storage,
- Extension::VK_KHR_shader_float16_int8,
- Extension::VK_KHR_storage_buffer_storage_class,
- Extension::VK_KHR_buffer_device_address,
- Extension::VK_KHR_variable_pointers,
- };
- return append_range(extensions, list);
+ return append_values(extensions, Extension::VK_KHR_16bit_storage,
+ Extension::VK_KHR_8bit_storage,
+ Extension::VK_KHR_shader_float16_int8,
+ Extension::VK_KHR_storage_buffer_storage_class,
+ Extension::VK_KHR_buffer_device_address,
+ Extension::VK_KHR_variable_pointers);
}
case TargetTripleArch::ARM_Valhall: {
// Example: https://vulkan.gpuinfo.org/displayreport.php?id=10312
- const Extension list[] = {
- Extension::VK_KHR_16bit_storage,
- Extension::VK_KHR_8bit_storage,
- Extension::VK_KHR_shader_float16_int8,
- Extension::VK_KHR_shader_integer_dot_product,
- Extension::VK_KHR_spirv_1_4,
- Extension::VK_KHR_storage_buffer_storage_class,
- Extension::VK_KHR_variable_pointers,
- };
- return append_range(extensions, list);
+ return append_values(extensions, Extension::VK_KHR_16bit_storage,
+ Extension::VK_KHR_8bit_storage,
+ Extension::VK_KHR_shader_float16_int8,
+ Extension::VK_KHR_shader_integer_dot_product,
+ Extension::VK_KHR_spirv_1_4,
+ Extension::VK_KHR_storage_buffer_storage_class,
+ Extension::VK_KHR_variable_pointers);
}
case TargetTripleArch::QC_Adreno: {
// Example: https://vulkan.gpuinfo.org/displayreport.php?id=10983 (11)
// Example: https://vulkan.gpuinfo.org/displayreport.php?id=16312 (12)
- const Extension list[] = {
- Extension::VK_KHR_16bit_storage,
- Extension::VK_KHR_shader_float16_int8,
- Extension::VK_KHR_spirv_1_4,
- Extension::VK_KHR_storage_buffer_storage_class,
- Extension::VK_KHR_variable_pointers,
- };
- append_range(extensions, list);
+ append_values(extensions, Extension::VK_KHR_16bit_storage,
+ Extension::VK_KHR_shader_float16_int8,
+ Extension::VK_KHR_spirv_1_4,
+ Extension::VK_KHR_storage_buffer_storage_class,
+ Extension::VK_KHR_variable_pointers);
if (triple.getOS() == TargetTripleOS::Android31) {
extensions.push_back(Extension::VK_KHR_8bit_storage);
}
@@ -179,18 +170,15 @@
return append_range(extensions, list);
}
- // Desktop GPUs typically support all extensions we care.
- const Extension desktop[] = {Extension::VK_KHR_16bit_storage,
- Extension::VK_KHR_8bit_storage,
- Extension::VK_KHR_shader_float16_int8,
- Extension::VK_KHR_shader_integer_dot_product,
- Extension::VK_KHR_spirv_1_4,
- Extension::VK_KHR_storage_buffer_storage_class,
- Extension::VK_KHR_buffer_device_address,
- Extension::VK_KHR_variable_pointers,
- Extension::VK_EXT_subgroup_size_control};
-
- llvm::append_range(extensions, desktop);
+ llvm::append_values(
+ extensions, // Desktop GPUs typically support all extensions we care.
+ Extension::VK_KHR_16bit_storage, Extension::VK_KHR_8bit_storage,
+ Extension::VK_KHR_shader_float16_int8,
+ Extension::VK_KHR_shader_integer_dot_product, Extension::VK_KHR_spirv_1_4,
+ Extension::VK_KHR_storage_buffer_storage_class,
+ Extension::VK_KHR_buffer_device_address,
+ Extension::VK_KHR_variable_pointers,
+ Extension::VK_EXT_subgroup_size_control);
if (getVendor(triple) == spirv::Vendor::NVIDIA ||
triple.getArch() == TargetTripleArch::AMD_RDNAv3) {
extensions.push_back(Extension::VK_KHR_cooperative_matrix);
diff --git a/compiler/src/iree/compiler/Modules/HAL/Inline/Transforms/InlineExecutables.cpp b/compiler/src/iree/compiler/Modules/HAL/Inline/Transforms/InlineExecutables.cpp
index f317326..89edd28 100644
--- a/compiler/src/iree/compiler/Modules/HAL/Inline/Transforms/InlineExecutables.cpp
+++ b/compiler/src/iree/compiler/Modules/HAL/Inline/Transforms/InlineExecutables.cpp
@@ -357,7 +357,7 @@
}
Value device = builder.create<IREE::Util::NullOp>(
loc, builder.getType<IREE::HAL::DeviceType>());
- auto workgroupCount =
+ std::array<Value, 3> workgroupCount =
exportOp.calculateWorkgroupCount(loc, device, workload, builder);
// For now we don't handle local memory.
@@ -385,13 +385,9 @@
}
int workgroupXYZOffset = workgroupArgs.size();
- workgroupArgs.push_back(nullptr); // workgroup_x, set below
- workgroupArgs.push_back(nullptr); // workgroup_y, set below
- workgroupArgs.push_back(nullptr); // workgroup_z, set below
- workgroupArgs.append(3, indexSet.get(1)); // workgroup_size_xyz
- workgroupArgs.push_back(workgroupCount[0]); // workgroup_count_x
- workgroupArgs.push_back(workgroupCount[1]); // workgroup_count_y
- workgroupArgs.push_back(workgroupCount[2]); // workgroup_count_z
+ workgroupArgs.append(3, nullptr); // workgroup_xyz, set below
+ workgroupArgs.append(3, indexSet.get(1)); // workgroup_size_xyz
+ llvm::append_range(workgroupArgs, workgroupCount); // workgroup_count_xyz
// Z -> Y -> Z loop nest.
builder.create<scf::ForOp>(
