Use interface ABI attrs in SPIR-V lowering. Use of the interface attributes provided by SPIR-V lowering in MLIR provides simplifies the code to lower the dispatch function to SPIR-V. The arguments of the dispatch function can just have attributes that describe the interface for the entry point function. These attributes are then materialized using a pass from SPIR-V lowering in MLIR. PiperOrigin-RevId: 285046031
diff --git a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/BUILD b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/BUILD index 699eb6e..1b4f485 100644 --- a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/BUILD +++ b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/BUILD
@@ -37,6 +37,7 @@ "@local_config_mlir//:IR", "@local_config_mlir//:Pass", "@local_config_mlir//:SPIRVDialect", + "@local_config_mlir//:SPIRVLowering", "@local_config_mlir//:SPIRVSerialization", "@local_config_mlir//:Support", "@org_tensorflow//tensorflow/compiler/mlir/xla:hlo",
diff --git a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.cpp b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.cpp index 02a5f89..b4a5b5d 100644 --- a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.cpp +++ b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.cpp
@@ -26,6 +26,7 @@ #include "iree/schemas/spirv_executable_def_generated.h" #include "llvm/ADT/STLExtras.h" #include "llvm/Support/CommandLine.h" +#include "mlir/Dialect/SPIRV/Passes.h" #include "mlir/Dialect/SPIRV/SPIRVOps.h" #include "mlir/Dialect/SPIRV/Serialization.h" #include "mlir/IR/Builders.h" @@ -143,6 +144,7 @@ conversionPassManager.addPass(xla_hlo::createLegalizeToStdPass()); conversionPassManager.addPass(createIndexComputationPass()); conversionPassManager.addPass(createIREEToSPIRVPass()); + conversionPassManager.addPass(spirv::createLowerABIAttributesPass()); conversionPassManager.addPass(createAdjustIntegerWidthPass()); if (failed(conversionPassManager.run(moduleOp))) { return moduleOp.emitError() << "failed to run conversion passes";
diff --git a/iree/compiler/Translation/SPIRV/BUILD b/iree/compiler/Translation/SPIRV/BUILD index c4a8dd0..3f0c0b6 100644 --- a/iree/compiler/Translation/SPIRV/BUILD +++ b/iree/compiler/Translation/SPIRV/BUILD
@@ -59,6 +59,7 @@ "@local_config_mlir//:Pass", "@local_config_mlir//:SPIRVDialect", "@local_config_mlir//:SPIRVDialectRegistration", + "@local_config_mlir//:SPIRVLowering", "@local_config_mlir//:SPIRVSerialization", "@local_config_mlir//:StandardDialectRegistration", "@local_config_mlir//:StandardOps",
diff --git a/iree/compiler/Translation/SPIRV/IREEToSPIRV.cpp b/iree/compiler/Translation/SPIRV/IREEToSPIRV.cpp index e50e38a..c4d28f2 100644 --- a/iree/compiler/Translation/SPIRV/IREEToSPIRV.cpp +++ b/iree/compiler/Translation/SPIRV/IREEToSPIRV.cpp
@@ -36,9 +36,8 @@ /// IREE::StoreOp needs to write to the spv.globalVariable created for the /// memref that holds the result of the dispatch function. LogicalResult IREEStoreOpSPIRVLowering::lowerOperation( - Operation *op, OpBuilder &builder, TensorIndexToScalarValueMap &valueCache, - DenseMap<Value *, spirv::GlobalVariableOp> &inputBuffers, - ArrayRef<spirv::GlobalVariableOp> outputBuffers) const { + Operation *op, OpBuilder &builder, + TensorIndexToScalarValueMap &valueCache) const { auto storeOp = cast<IREE::StoreOutputOp>(op); auto src = storeOp.src(); SmallVector<AffineMap, 1> indices; @@ -48,10 +47,10 @@ "expected to compute a single element of the tensor that is stored " "into the output memref"); } - auto var = inputBuffers.lookup(storeOp.dst()); + auto var = valueCache.getBufferForArgument(storeOp.dst()); if (!var) { return storeOp.emitError( - "unable to find spv.globalVariable that corresponds to the dst memref"); + "unable to find buffer that corresponds to the dst memref"); } auto ptr = genPointerOffset(builder, storeOp.getLoc(), valueCache, indices[0], var); @@ -65,9 +64,8 @@ /// IREE::ReturnOp in dispatch functions lowered to SPIR-V should have no /// operands. LogicalResult IREEReturnOpSPIRVLowering::lowerOperation( - Operation *op, OpBuilder &builder, TensorIndexToScalarValueMap &valueCache, - DenseMap<Value *, spirv::GlobalVariableOp> &inputBuffers, - ArrayRef<spirv::GlobalVariableOp> outputBuffers) const { + Operation *op, OpBuilder &builder, + TensorIndexToScalarValueMap &valueCache) const { return success(); }
diff --git a/iree/compiler/Translation/SPIRV/IREEToSPIRV.h b/iree/compiler/Translation/SPIRV/IREEToSPIRV.h index ffb8013..ade0baa 100644 --- a/iree/compiler/Translation/SPIRV/IREEToSPIRV.h +++ b/iree/compiler/Translation/SPIRV/IREEToSPIRV.h
@@ -46,9 +46,7 @@ LogicalResult lowerOperation( Operation *op, OpBuilder &builder, - TensorIndexToScalarValueMap &valueCache, - DenseMap<Value *, spirv::GlobalVariableOp> &inputBuffers, - ArrayRef<spirv::GlobalVariableOp> outputBuffers) const override; + TensorIndexToScalarValueMap &valueCache) const override; }; /// Translation of iree.store_output operation. @@ -59,9 +57,7 @@ LogicalResult lowerOperation( Operation *op, OpBuilder &builder, - TensorIndexToScalarValueMap &valueCache, - DenseMap<Value *, spirv::GlobalVariableOp> &inputBuffers, - ArrayRef<spirv::GlobalVariableOp> outputBuffers) const override; + TensorIndexToScalarValueMap &valueCache) const override; }; } // namespace iree_compiler
diff --git a/iree/compiler/Translation/SPIRV/SPIRVExecutableTranslation.cpp b/iree/compiler/Translation/SPIRV/SPIRVExecutableTranslation.cpp index 575816a..7b46e5b 100644 --- a/iree/compiler/Translation/SPIRV/SPIRVExecutableTranslation.cpp +++ b/iree/compiler/Translation/SPIRV/SPIRVExecutableTranslation.cpp
@@ -30,6 +30,7 @@ #include "llvm/ADT/StringRef.h" #include "llvm/Support/Debug.h" #include "llvm/Support/ErrorHandling.h" +#include "mlir/Dialect/SPIRV/Passes.h" #include "mlir/Dialect/SPIRV/SPIRVOps.h" #include "mlir/Dialect/SPIRV/Serialization.h" #include "mlir/IR/Attributes.h" @@ -206,6 +207,7 @@ spirvGenPasses->addPass(xla_hlo::createLegalizeToStdPass()); spirvGenPasses->addPass(createIndexComputationPass()); spirvGenPasses->addPass(createIREEToSPIRVPass()); + spirvGenPasses->addPass(spirv::createLowerABIAttributesPass()); spirvGenPasses->addPass(createAdjustIntegerWidthPass()); if (failed(runPassPipeline(options(), spirvGenPasses.get(), module))) { executableOp.emitError() << "Failed to generate spv.module";
diff --git a/iree/compiler/Translation/SPIRV/SPIRVLowering.cpp b/iree/compiler/Translation/SPIRV/SPIRVLowering.cpp index 4331dbd..80d5463 100644 --- a/iree/compiler/Translation/SPIRV/SPIRVLowering.cpp +++ b/iree/compiler/Translation/SPIRV/SPIRVLowering.cpp
@@ -19,6 +19,8 @@ //===----------------------------------------------------------------------===// #include "iree/compiler/Translation/SPIRV/SPIRVLowering.h" +#include "mlir/Dialect/SPIRV/SPIRVLowering.h" + namespace mlir { namespace iree_compiler { @@ -28,12 +30,10 @@ Value *genPointerOffset(OpBuilder &builder, Location loc, TensorIndexToScalarValueMap &valueCache, - AffineMap indexMap, spirv::GlobalVariableOp var) { - auto basePtr = builder.create<spirv::AddressOfOp>(loc, var); - auto varPtrType = var.type().cast<spirv::PointerType>().getPointeeType(); + AffineMap indexMap, Value *buffer) { + auto varPtrType = + buffer->getType().cast<spirv::PointerType>().getPointeeType(); // The variable has to be a struct type with a single element. - assert(varPtrType.isa<spirv::StructType>() && - "expected variable type to be a spv.ptr<spv.struct<...>>"); auto varStructType = varPtrType.cast<spirv::StructType>(); assert(varStructType.getNumElements() == 1 && "expected variable type to be a spv.ptr of spv.struct with a single " @@ -54,7 +54,7 @@ accessIndices.push_back(valueCache.getAffineExprValue( builder.saveInsertionPoint(), loc, indexExpr)); } - return builder.create<spirv::AccessChainOp>(loc, basePtr, accessIndices); + return builder.create<spirv::AccessChainOp>(loc, buffer, accessIndices); } /// Returns the type of the global variable to use for an argument of the @@ -88,27 +88,6 @@ spirv::StorageClass::StorageBuffer); } -/// Returns a spirv::GlobalVariable op for a give argument of type `argType` and -/// argument idnex `argIndex`. -static spirv::GlobalVariableOp createGlobalVariableForArg(Location loc, - OpBuilder &builder, - unsigned argIndex, - Type argType, - StringRef varName) { - auto argMemrefType = argType.dyn_cast<MemRefType>(); - if (!argMemrefType) { - emitError(loc, "unhandled non-memref type argument in SPIR-V lowering"); - return nullptr; - } - auto varType = convertArgTypeToVariableType(loc, argMemrefType); - if (!varType) { - return nullptr; - } - auto var = builder.create<spirv::GlobalVariableOp>(loc, varType, varName, 0, - argIndex); - return var; -} - static spirv::GlobalVariableOp createGlobalInvocationID(OpBuilder &builder, Location loc) { auto i32Type = builder.getIntegerType(32); @@ -123,48 +102,63 @@ OpBuilder &builder, FuncOp &fn, TensorIndexToScalarValueMap &valueCache) { auto loc = fn.getLoc(); - // Convert functions arguments and return values to - // spirv::GlobalVariables. All global variables are given a descriptor set - // of 0 and binding is the argument number. - auto fnType = fn.getType(); - for (auto argType : enumerate(fnType.getInputs())) { - auto &var = inputArgToVariable[fn.getArgument(argType.index())]; - auto varName = - fn.getName().str() + "_arg_" + std::to_string(argType.index()); - var = createGlobalVariableForArg(loc, builder, argType.index(), - argType.value(), varName); - if (!var) { - return failure(); - } - } - resultIndexToVariable.resize(fnType.getNumResults()); - for (auto resType : enumerate(fnType.getResults())) { - auto &var = resultIndexToVariable[resType.index()]; - auto varName = - fn.getName().str() + "_result_" + std::to_string(resType.index()); - var = createGlobalVariableForArg(loc, builder, - resType.index() + fnType.getNumInputs(), - resType.value(), varName); - if (!var) { - return failure(); - } - } - // Create the Global invocation ID. auto globalInvocationID = createGlobalInvocationID(builder, fn.getLoc()); interface.push_back(builder.getSymbolRefAttr(globalInvocationID.sym_name())); - auto entryFnType = - builder.getFunctionType(ArrayRef<Type>(), ArrayRef<Type>()); + // Add ABI attributes to function arguments and return values. + auto fnType = fn.getType(); + SmallVector<Type, 2> entryFnArgTypes; + SmallVector<spirv::InterfaceVarABIAttr, 2> entryFnArgAttrs; + entryFnArgTypes.reserve(fnType.getNumInputs()); + entryFnArgAttrs.reserve(fnType.getNumInputs()); + for (auto argType : enumerate(fnType.getInputs())) { + auto argMemRefType = argType.value().dyn_cast<MemRefType>(); + if (!argMemRefType) { + return fn.emitError("expected arguments to be of memref types"); + } + auto convertedArgType = convertArgTypeToVariableType(loc, argMemRefType); + if (!convertedArgType) { + return failure(); + } + entryFnArgTypes.emplace_back(convertedArgType); + entryFnArgAttrs.emplace_back(spirv::getInterfaceVarABIAttr( + 0, argType.index(), spirv::StorageClass::StorageBuffer, + builder.getContext())); + } + + // TODO(ravishankarm) : Handle return types. The SPIR-V ABI attribute lowering + // needs to support attributes on return values. + if (fnType.getNumResults()) { + return fn.emitError("unimplemented handling for return types"); + } + + auto entryFnType = builder.getFunctionType(entryFnArgTypes, ArrayRef<Type>()); auto entryFn = builder.create<FuncOp>(loc, fn.getName(), entryFnType, ArrayRef<NamedAttribute>()); + entryFn.addEntryBlock(); + + SmallVector<int32_t, 3> workGroupSize; + if (failed(getWorkGroupSize(fn, workGroupSize))) { + return failure(); + } + auto entryFnAttr = + spirv::getEntryPointABIAttr(workGroupSize, builder.getContext()); + if (failed(setABIAttrs(entryFn, entryFnAttr, entryFnArgAttrs))) { + return failure(); + } + // Set the argument to buffer mapping + for (auto arg : enumerate(fn.getArguments())) { + valueCache.setBufferForArgument(arg.value(), + entryFn.getArgument(arg.index())); + } // Start a scope to create an insertion guard to reset the builder once the // function is lowered. { assert(globalInvocationIDs.empty()); OpBuilder::InsertionGuard funcInsertGuard(builder); - builder.setInsertionPointToStart(entryFn.addEntryBlock()); + builder.setInsertionPointToStart(&entryFn.getBody().front()); auto globalInvocationIDPtr = builder.create<spirv::AddressOfOp>(loc, globalInvocationID); @@ -187,26 +181,6 @@ } } - // Create the entry point instructions for the entry function. - if (failed(createEntryPoint(fn, builder, loc, entryFn))) { - return failure(); - } - return success(); -} - -LogicalResult SPIRVCodegenImpl::createEntryPoint(FuncOp fn, OpBuilder &builder, - Location loc, FuncOp entryFn) { - builder.create<spirv::EntryPointOp>(loc, spirv::ExecutionModel::GLCompute, - entryFn, interface); - - // Get the workgroup size. - SmallVector<int32_t, 3> workGroupSize; - if (failed(getWorkGroupSize(fn, workGroupSize))) { - return failure(); - } - builder.create<spirv::ExecutionModeOp>( - loc, entryFn, spirv::ExecutionMode::LocalSize, workGroupSize); - interface.clear(); return success(); } @@ -301,9 +275,9 @@ if (val) { return val; } - auto var = inputArgToVariable.lookup(origArg); + auto var = valueCache.getBufferForArgument(origArg); if (!var) { - emitError(loc, "undefined SPIR-V global variable for tensor argument"); + emitError(loc, "unable to find buffer for tensor argument"); return nullptr; } auto ptr = genPointerOffset(builder, loc, valueCache, indexMap, var); @@ -517,35 +491,5 @@ return success(); } -//===----------------------------------------------------------------------===// -// ReturnOp -//===----------------------------------------------------------------------===// -LogicalResult ReturnOpSPIRVLowering::lowerOperation( - Operation *op, OpBuilder &builder, TensorIndexToScalarValueMap &valueCache, - DenseMap<Value *, spirv::GlobalVariableOp> &inputBuffers, - ArrayRef<spirv::GlobalVariableOp> outputBuffers) const { - auto returnOp = cast<ReturnOp>(op); - if (returnOp.getNumOperands() != 1) { - return returnOp.emitError( - "unhandled lowering of return statement with multiple returns"); - } - auto returnTensor = returnOp.getOperand(0); - SmallVector<AffineMap, 1> indices; - index_computation_attribute::getIndexMapsForValue(returnTensor, indices); - if (indices.size() != 1) { - return returnOp.emitError( - "expected to compute a single element of the return tensor"); - } - assert(outputBuffers.size() == 1 && "Expected a single output buffer"); - auto var = outputBuffers[0]; - auto ptr = - genPointerOffset(builder, returnOp.getLoc(), valueCache, indices[0], var); - auto scalarVal = valueCache.getValueAtIndex(returnTensor, indices[0]); - builder.create<spirv::StoreOp>(returnOp.getLoc(), ptr, scalarVal, - /*memory_access = */ nullptr, - /*alignment = */ nullptr); - builder.create<spirv::ReturnOp>(returnOp.getLoc()); - return success(); -} } // namespace iree_compiler } // namespace mlir
diff --git a/iree/compiler/Translation/SPIRV/SPIRVLowering.h b/iree/compiler/Translation/SPIRV/SPIRVLowering.h index ee1f57a..2224d0c 100644 --- a/iree/compiler/Translation/SPIRV/SPIRVLowering.h +++ b/iree/compiler/Translation/SPIRV/SPIRVLowering.h
@@ -55,9 +55,7 @@ /// operations. virtual LogicalResult lowerOperation( Operation *op, OpBuilder &builder, - TensorIndexToScalarValueMap &valueCache, - DenseMap<Value *, spirv::GlobalVariableOp> &inputBuffers, - ArrayRef<spirv::GlobalVariableOp> outputBuffers) const { + TensorIndexToScalarValueMap &valueCache) const { return failure(); } }; @@ -233,19 +231,7 @@ /// location of a spv.globalVariable. Value *genPointerOffset(OpBuilder &builder, Location loc, TensorIndexToScalarValueMap &valueCache, - AffineMap indexMap, spirv::GlobalVariableOp var); - -/// Lower return statements during SPIR-V codegeneration. -class ReturnOpSPIRVLowering : public SPIRVOpLowering<ReturnOp> { - public: - using SPIRVOpLowering<ReturnOp>::SPIRVOpLowering; - - LogicalResult lowerOperation( - Operation *op, OpBuilder &builder, - TensorIndexToScalarValueMap &valueCache, - DenseMap<Value *, spirv::GlobalVariableOp> &inputBuffers, - ArrayRef<spirv::GlobalVariableOp> outputBuffers) const override; -}; + AffineMap indexMap, Value *buffer); namespace detail { /// Implementation class for generating SPIR-V kernel. @@ -353,8 +339,7 @@ // Zero return case. if (!op->getNumResults()) { - return opCodegenList[opName]->lowerOperation( - op, builder, valueCache, inputArgToVariable, resultIndexToVariable); + return opCodegenList[opName]->lowerOperation(op, builder, valueCache); } // Single return case.
diff --git a/iree/compiler/Translation/SPIRV/TensorIndexToScalarValueMap.h b/iree/compiler/Translation/SPIRV/TensorIndexToScalarValueMap.h index d542629..c51cb32 100644 --- a/iree/compiler/Translation/SPIRV/TensorIndexToScalarValueMap.h +++ b/iree/compiler/Translation/SPIRV/TensorIndexToScalarValueMap.h
@@ -36,6 +36,9 @@ explicit TensorIndexToScalarValueMap(MLIRContext *context) : affineExprCodegen(context) {} + /// Returns the buffer associated with an argument in the dispatch function. + Value *getBufferForArgument(Value *arg) { return argToBufferMap.lookup(arg); } + /// Returns the value in the lowered function that represents the scalar value /// of the `tensor` in the original function at a given `index` Value *getValueAtIndex(Value *tensor, AffineMap index) { @@ -58,6 +61,11 @@ return affineExprCodegen.getExprValue(ip, loc, expr); } + /// Records the `buffer` to use for an `argument` in the dispatch function. + void setBufferForArgument(Value *argument, Value *buffer) { + argToBufferMap[argument] = buffer; + } + /// Records the `scalar` value in the lowered function that represents the /// value of the `tensor` in the original function at a given `index` void setValueAtIndex(Value *tensor, AffineMap index, Value *scalar) { @@ -164,6 +172,7 @@ AffineExprCodegen affineExprCodegen; DenseMap<Value *, DenseMap<AffineMap, Value *>> tensorIndexToScalarValueMap; + DenseMap<Value *, Value *> argToBufferMap; }; } // namespace iree_compiler
diff --git a/iree/compiler/Translation/SPIRV/test/broadcast.mlir b/iree/compiler/Translation/SPIRV/test/broadcast.mlir index d6cdd0c..e8c67b7 100644 --- a/iree/compiler/Translation/SPIRV/test/broadcast.mlir +++ b/iree/compiler/Translation/SPIRV/test/broadcast.mlir
@@ -17,9 +17,9 @@ module { // CHECK:spv.module "Logical" "GLSL450" // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) - // CHECK-DAG: spv.globalVariable [[ARG1VAR:@.*]] bind(0, 1) - // CHECK: func [[FN:@broadcast_2D_3D]]() + // CHECK: func [[FN:@broadcast_2D_3D]] + // CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<42 x i32 [4]> [168]> [0]>, StorageBuffer> + // CHECK-SAME: [[ARG1:%.*]]: !spv.ptr<!spv.struct<!spv.array<3 x !spv.array<12 x !spv.array<42 x i32 [4]> [168]> [2016]> [0]>, StorageBuffer> func @broadcast_2D_3D(%arg0: memref<12x42xi32>, %arg1: memref<3x12x42xi32>) attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 3]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[GLOBALIDPTR:%.*]] = spv._address_of [[GLOBALIDVAR]] @@ -28,15 +28,13 @@ // CHECK: [[GLOBALIDY:%.*]] = spv.CompositeExtract [[GLOBALID]][1 : i32] // CHECK: [[GLOBALIDZ:%.*]] = spv.CompositeExtract [[GLOBALID]][2 : i32] // CHECK: spv.selection - // CHECK: [[ARG0PTR:%.*]] = spv._address_of [[ARG0VAR]] // CHECK: [[ZERO1:%.*]] = spv.constant 0 : i32 - // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0PTR]]{{\[}}[[ZERO1]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} + // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]{{\[}}[[ZERO1]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} // CHECK: [[VAL:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]] %0 = iree.load_input(%arg0 : memref<12x42xi32>) : tensor<12x42xi32> %1 = "xla_hlo.broadcast"(%0) {broadcast_sizes = dense<[3]> : tensor<1xi64>} : (tensor<12x42xi32>) -> tensor<3x12x42xi32> - // CHECK: [[ARG1PTR:%.*]] = spv._address_of [[ARG1VAR]] // CHECK: [[ZERO2:%.*]] = spv.constant 0 : i32 - // CHECK: [[ARG1STOREPTR:%.*]] = spv.AccessChain [[ARG1PTR]]{{\[}}[[ZERO2]], [[GLOBALIDZ]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} + // CHECK: [[ARG1STOREPTR:%.*]] = spv.AccessChain [[ARG1]]{{\[}}[[ZERO2]], [[GLOBALIDZ]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} // CHECK: spv.Store "StorageBuffer" [[ARG1STOREPTR]], [[VAL]] iree.store_output(%1 : tensor<3x12x42xi32>, %arg1 : memref<3x12x42xi32>) iree.return @@ -48,9 +46,9 @@ module { // CHECK:spv.module "Logical" "GLSL450" // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer> - // CHECK-DAG: spv.globalVariable [[ARG1VAR:@.*]] bind(0, 1) - // CHECK: func [[FN:@broadcast_scalar_3D]]() + // CHECK: func [[FN:@broadcast_scalar_3D]] + // CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer> + // CHECK-SAME: [[ARG1:%.*]]: !spv.ptr<!spv.struct<!spv.array<3 x !spv.array<12 x !spv.array<42 x i32 [4]> [168]> [2016]> [0]>, StorageBuffer> func @broadcast_scalar_3D(%arg0: memref<i32>, %arg1: memref<3x12x42xi32>) attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 3]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[GLOBALIDPTR:%.*]] = spv._address_of [[GLOBALIDVAR]] @@ -59,15 +57,13 @@ // CHECK: [[GLOBALIDY:%.*]] = spv.CompositeExtract [[GLOBALID]][1 : i32] // CHECK: [[GLOBALIDZ:%.*]] = spv.CompositeExtract [[GLOBALID]][2 : i32] // CHECK: spv.selection - // CHECK: [[ARG0PTR:%.*]] = spv._address_of [[ARG0VAR]] // CHECK: [[ZERO:%.*]] = spv.constant 0 : i32 - // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0PTR]]{{\[}}[[ZERO]]{{\]}} + // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]{{\[}}[[ZERO]]{{\]}} // CHECK: [[VAL:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]] %0 = iree.load_input(%arg0 : memref<i32>) : tensor<i32> %1 = "xla_hlo.broadcast"(%0) {broadcast_sizes = dense<[3, 12, 42]>: tensor<3xi64>} : (tensor<i32>) -> tensor<3x12x42xi32> - // CHECK: [[ARG1PTR:%.*]] = spv._address_of [[ARG1VAR]] // CHECK: [[ZERO1:%.*]] = spv.constant 0 : i32 - // CHECK: [[ARG1STOREPTR:%.*]] = spv.AccessChain [[ARG1PTR]]{{\[}}[[ZERO1]], [[GLOBALIDZ]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} + // CHECK: [[ARG1STOREPTR:%.*]] = spv.AccessChain [[ARG1]]{{\[}}[[ZERO1]], [[GLOBALIDZ]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} // CHECK: spv.Store "StorageBuffer" [[ARG1STOREPTR]], [[VAL]] iree.store_output(%1 : tensor<3x12x42xi32>, %arg1 : memref<3x12x42xi32>) iree.return
diff --git a/iree/compiler/Translation/SPIRV/test/broadcast_in_dim.mlir b/iree/compiler/Translation/SPIRV/test/broadcast_in_dim.mlir index d3e5416..4466607 100644 --- a/iree/compiler/Translation/SPIRV/test/broadcast_in_dim.mlir +++ b/iree/compiler/Translation/SPIRV/test/broadcast_in_dim.mlir
@@ -17,9 +17,9 @@ module { // CHECK:spv.module "Logical" "GLSL450" // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) - // CHECK-DAG: spv.globalVariable [[ARG1VAR:@.*]] bind(0, 1) - // CHECK: func [[FN:@broadcast_in_dim_2D_3D]]() + // CHECK: func [[FN:@broadcast_in_dim_2D_3D]] + // CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<42 x i32 [4]> [168]> [0]>, StorageBuffer> + // CHECK-SAME: [[ARG1:%.*]]: !spv.ptr<!spv.struct<!spv.array<3 x !spv.array<12 x !spv.array<42 x i32 [4]> [168]> [2016]> [0]>, StorageBuffer> func @broadcast_in_dim_2D_3D(%arg0: memref<12x42xi32>, %arg1: memref<3x12x42xi32>) attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 3]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[GLOBALIDPTR:%.*]] = spv._address_of [[GLOBALIDVAR]] @@ -28,15 +28,13 @@ // CHECK: [[GLOBALIDY:%.*]] = spv.CompositeExtract [[GLOBALID]][1 : i32] // CHECK: [[GLOBALIDZ:%.*]] = spv.CompositeExtract [[GLOBALID]][2 : i32] // CHECK: spv.selection - // CHECK: [[ARG0PTR:%.*]] = spv._address_of [[ARG0VAR]] // CHECK: [[ZERO1:%.*]] = spv.constant 0 : i32 - // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0PTR]]{{\[}}[[ZERO1]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} + // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]{{\[}}[[ZERO1]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} // CHECK: [[VAL:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]] %0 = iree.load_input(%arg0 : memref<12x42xi32>) : tensor<12x42xi32> %1 = "xla_hlo.broadcast_in_dim"(%0) {broadcast_dimensions = dense<[1, 2]> : tensor<2xi64>} : (tensor<12x42xi32>) -> tensor<3x12x42xi32> - // CHECK: [[ARG1PTR:%.*]] = spv._address_of [[ARG1VAR]] // CHECK: [[ZERO2:%.*]] = spv.constant 0 : i32 - // CHECK: [[ARG1STOREPTR:%.*]] = spv.AccessChain [[ARG1PTR]]{{\[}}[[ZERO2]], [[GLOBALIDZ]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} + // CHECK: [[ARG1STOREPTR:%.*]] = spv.AccessChain [[ARG1]]{{\[}}[[ZERO2]], [[GLOBALIDZ]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} // CHECK: spv.Store "StorageBuffer" [[ARG1STOREPTR]], [[VAL]] iree.store_output(%1 : tensor<3x12x42xi32>, %arg1 : memref<3x12x42xi32>) iree.return @@ -48,9 +46,9 @@ module { // CHECK:spv.module "Logical" "GLSL450" // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer> - // CHECK-DAG: spv.globalVariable [[ARG1VAR:@.*]] bind(0, 1) - // CHECK: func [[FN:@broadcast_in_dim_scalar_3D]]() + // CHECK: func [[FN:@broadcast_in_dim_scalar_3D]] + // CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer> + // CHECK-SAME: [[ARG1:%.*]]: !spv.ptr<!spv.struct<!spv.array<3 x !spv.array<12 x !spv.array<42 x i32 [4]> [168]> [2016]> [0]>, StorageBuffer> func @broadcast_in_dim_scalar_3D(%arg0: memref<i32>, %arg1: memref<3x12x42xi32>) attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 3]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[GLOBALIDPTR:%.*]] = spv._address_of [[GLOBALIDVAR]] @@ -59,15 +57,13 @@ // CHECK: [[GLOBALIDY:%.*]] = spv.CompositeExtract [[GLOBALID]][1 : i32] // CHECK: [[GLOBALIDZ:%.*]] = spv.CompositeExtract [[GLOBALID]][2 : i32] // CHECK: spv.selection - // CHECK: [[ARG0PTR:%.*]] = spv._address_of [[ARG0VAR]] // CHECK: [[ZERO:%.*]] = spv.constant 0 : i32 - // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0PTR]]{{\[}}[[ZERO]]{{\]}} + // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]{{\[}}[[ZERO]]{{\]}} // CHECK: [[VAL:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]] %0 = iree.load_input(%arg0 : memref<i32>) : tensor<i32> %1 = "xla_hlo.broadcast_in_dim"(%0) : (tensor<i32>) -> tensor<3x12x42xi32> - // CHECK: [[ARG1PTR:%.*]] = spv._address_of [[ARG1VAR]] // CHECK: [[ZERO1:%.*]] = spv.constant 0 : i32 - // CHECK: [[ARG1STOREPTR:%.*]] = spv.AccessChain [[ARG1PTR]]{{\[}}[[ZERO1]], [[GLOBALIDZ]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} + // CHECK: [[ARG1STOREPTR:%.*]] = spv.AccessChain [[ARG1]]{{\[}}[[ZERO1]], [[GLOBALIDZ]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} // CHECK: spv.Store "StorageBuffer" [[ARG1STOREPTR]], [[VAL]] iree.store_output(%1 : tensor<3x12x42xi32>, %arg1 : memref<3x12x42xi32>) iree.return
diff --git a/iree/compiler/Translation/SPIRV/test/concatenate.mlir b/iree/compiler/Translation/SPIRV/test/concatenate.mlir index 8eba18b..207fe85 100644 --- a/iree/compiler/Translation/SPIRV/test/concatenate.mlir +++ b/iree/compiler/Translation/SPIRV/test/concatenate.mlir
@@ -16,24 +16,23 @@ module { // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) - // CHECK-DAG: spv.globalVariable [[ARG1VAR:@.*]] bind(0, 1) + // CHECK: func @concatenate + // CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<1 x !spv.array<64 x f32 [4]> [256]> [0]>, StorageBuffer> + // CHECK-SAME: [[ARG1:%.*]]: !spv.ptr<!spv.struct<!spv.array<1 x !spv.array<10 x f32 [4]> [40]> [0]>, StorageBuffer> func @concatenate(%arg0: memref<1x64xf32>, %arg1 : memref<1x10xf32>, %arg2 : memref<1x74xf32>) attributes {iree.executable.export, iree.executable.workload = dense<[1, 74]> : tensor<2xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[GLOBALIDPTR:%.*]] = spv._address_of [[GLOBALIDVAR]] // CHECK: [[GLOBALID:%.*]] = spv.Load "Input" [[GLOBALIDPTR]] // CHECK: [[GLOBALIDX:%.*]] = spv.CompositeExtract [[GLOBALID]][0 : i32] // CHECK: [[GLOBALIDY:%.*]] = spv.CompositeExtract [[GLOBALID]][1 : i32] - // CHECK: [[ARG0PTR:%.*]] = spv._address_of [[ARG0VAR]] // CHECK: [[ZERO1:%.*]] = spv.constant 0 : i32 // CHECK: [[ZERO2:%.*]] = spv.constant 0 : i32 - // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0PTR]]{{\[}}[[ZERO1]], [[ZERO2]], [[GLOBALIDY]]{{\]}} + // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]{{\[}}[[ZERO1]], [[ZERO2]], [[GLOBALIDY]]{{\]}} // CHECK: [[INPUTVAL0:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]] : f32 - // CHECK: [[ARG1PTR:%.*]] = spv._address_of [[ARG1VAR]] // CHECK: [[ZERO3:%.*]] = spv.constant 0 : i32 // CHECK: [[NEGATIVE_SIXTY_FOUR:%.*]] = spv.constant -64 : i32 // CHECK: [[VAR1:%.*]] = spv.IAdd [[GLOBALIDY]], [[NEGATIVE_SIXTY_FOUR]] : i32 - // CHECK: [[ARG1LOADPTR:%.*]] = spv.AccessChain [[ARG1PTR]]{{\[}}[[ZERO3]], [[ZERO2]], [[VAR1]]{{\]}} + // CHECK: [[ARG1LOADPTR:%.*]] = spv.AccessChain [[ARG1]]{{\[}}[[ZERO3]], [[ZERO2]], [[VAR1]]{{\]}} // CHECK: [[INPUTVAL1:%.*]] = spv.Load "StorageBuffer" [[ARG1LOADPTR]] : f32 // CHECK: [[TRUE:%.*]] = spv.constant true // CHECK: [[SIXTY_FOUR:%.*]] = spv.constant 64 : i32
diff --git a/iree/compiler/Translation/SPIRV/test/constant.mlir b/iree/compiler/Translation/SPIRV/test/constant.mlir index 67665f8..ce1f7a5 100644 --- a/iree/compiler/Translation/SPIRV/test/constant.mlir +++ b/iree/compiler/Translation/SPIRV/test/constant.mlir
@@ -15,9 +15,6 @@ // RUN: iree-opt -xla-legalize-to-std -split-input-file -iree-index-computation -simplify-spirv-affine-exprs=false -convert-iree-to-spirv -verify-diagnostics -o - %s | IreeFileCheck %s module { - // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) - // CHECK-DAG: spv.globalVariable [[ARG1VAR:@.*]] bind(0, 1) func @const_f32(%arg0: memref<2x3xf32>, %arg1: memref<2x3xf32>) attributes {iree.executable.export, iree.executable.workload = dense<[2, 3]> : tensor<2xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[NUM0:%.*]] = spv.Load "StorageBuffer" {{%.*}} : f32 @@ -37,9 +34,6 @@ // ----- module { - // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) - // CHECK-DAG: spv.globalVariable [[ARG1VAR:@.*]] bind(0, 1) func @splat_const_f32(%arg0: memref<2x3xf32>, %arg1: memref<2x3xf32>) attributes {iree.executable.export, iree.executable.workload = dense<[2, 3]> : tensor<2xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[NUM0:%.*]] = spv.Load "StorageBuffer" {{%.*}} : f32 @@ -56,8 +50,6 @@ // ----- module { - // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) func @const_i32(%arg0: memref<2x3xi32>, %arg1: memref<2x3xi32>) attributes {iree.executable.export, iree.executable.workload = dense<[2, 3]> : tensor<2xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[NUM0:%.*]] = spv.Load "StorageBuffer" {{%.*}} : i32 @@ -79,7 +71,6 @@ module { // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) func @splat_const_i32(%arg0: memref<2x3xi32>, %arg1: memref<2x3xi32>) attributes {iree.executable.export, iree.executable.workload = dense<[2, 3]> : tensor<2xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[NUM0:%.*]] = spv.Load "StorageBuffer" {{%.*}} : i32
diff --git a/iree/compiler/Translation/SPIRV/test/copy.mlir b/iree/compiler/Translation/SPIRV/test/copy.mlir index fb20cd2..eb1fb6c 100644 --- a/iree/compiler/Translation/SPIRV/test/copy.mlir +++ b/iree/compiler/Translation/SPIRV/test/copy.mlir
@@ -17,9 +17,10 @@ module { // CHECK:spv.module "Logical" "GLSL450" // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<42 x i32 [4]> [168]> [0]>, StorageBuffer> - // CHECK-DAG: spv.globalVariable [[ARG1VAR:@.*]] bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<42 x i32 [4]> [168]> [0]>, StorageBuffer> - // CHECK: func [[FN:@simple_load_store]]() + // CHECK: func [[FN:@simple_load_store]] + // CHECK-SAME: [[ARG0:%[a-zA-Z0-9]*]]: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<42 x i32 [4]> [168]> [0]>, StorageBuffer> + // CHECK-SAME: [[ARG1:%[a-zA-Z0-9]*]]: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<42 x i32 [4]> [168]> [0]>, StorageBuffer> + // CHECK-SAME: spirv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>} func @simple_load_store(%arg0: memref<12x42xi32>, %arg1: memref<12x42xi32>) attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[GLOBALIDPTR:%.*]] = spv._address_of [[GLOBALIDVAR]] @@ -36,23 +37,19 @@ // CHECK: spv.selection { // CHECK: spv.BranchConditional [[COND2]], [[BB1:\^.*]], [[BB2:\^.*]] // CHECK: [[BB1]]: - // CHECK: [[ARG0PTR:%.*]] = spv._address_of [[ARG0VAR]] // CHECK: [[ZERO1:%.*]] = spv.constant 0 : i32 - // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0PTR]]{{\[}}[[ZERO1]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} + // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]{{\[}}[[ZERO1]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} // CHECK: [[VAL:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]] %0 = iree.load_input(%arg0 : memref<12x42xi32>) : tensor<12x42xi32> %1 = "xla_hlo.copy"(%0) : (tensor<12x42xi32>) -> tensor<12x42xi32> - // CHECK: [[ARG1PTR:%.*]] = spv._address_of [[ARG1VAR]] // CHECK: [[ZERO2:%.*]] = spv.constant 0 : i32 - // CHECK: [[ARG1STOREPTR:%.*]] = spv.AccessChain [[ARG1PTR]]{{\[}}[[ZERO2]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} + // CHECK: [[ARG1STOREPTR:%.*]] = spv.AccessChain [[ARG1]]{{\[}}[[ZERO2]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} // CHECK: spv.Store "StorageBuffer" [[ARG1STOREPTR]], [[VAL]] // CHECK: spv.Branch [[BB2]] iree.store_output(%1 : tensor<12x42xi32>, %arg1 : memref<12x42xi32>) // CHECK: spv.Return iree.return } - // CHECK: spv.EntryPoint "GLCompute" [[FN]], [[GLOBALIDVAR]] - // CHECK: spv.ExecutionMode [[FN]] "LocalSize", 32, 1, 1 } // -----
diff --git a/iree/compiler/Translation/SPIRV/test/extract_element.mlir b/iree/compiler/Translation/SPIRV/test/extract_element.mlir index 8493564..b519142 100644 --- a/iree/compiler/Translation/SPIRV/test/extract_element.mlir +++ b/iree/compiler/Translation/SPIRV/test/extract_element.mlir
@@ -15,16 +15,15 @@ // RUN: iree-opt -split-input-file -iree-index-computation -simplify-spirv-affine-exprs=false -convert-iree-to-spirv -verify-diagnostics -o - %s | IreeFileCheck %s module { - // CHECK: spv.globalVariable [[VAR0:@.*]] bind(0, 0) - // CHECK: spv.globalVariable [[VAR1:@.*]] bind(0, 1) + // CHECK-LABEL: func @extract_element + // CHECK-SAME: [[ARG0:%[a-zA-Z0-9]*]]: !spv.ptr<!spv.struct<i1 [0]>, StorageBuffer> + // CHECK-SAME: [[ARG1:%[a-zA-Z0-9]*]]: !spv.ptr<!spv.struct<i1 [0]>, StorageBuffer> func @extract_element(%arg0: memref<i1>, %arg1: memref<i1>) attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi64>, iree.executable.workload = dense<1> : tensor<3xi32>, iree.num_dims = 3 : i32, iree.ordinal = 0 : i32} { %0 = "iree.load_input"(%arg0) : (memref<i1>) -> tensor<i1> - // CHECK: [[ARG0:%.*]] = spv._address_of [[VAR0]] // CHECK: [[ZERO:%.*]] = spv.constant 0 : i32 // CHECK: {{%.*}} = spv.AccessChain [[ARG0]]{{\[}}[[ZERO]]{{\]}} %1 = "std.extract_element"(%0) : (tensor<i1>) -> i1 - // CHECK: [[ARG1:%.*]] = spv._address_of [[VAR1]] // CHECK: {{%.*}} = spv.AccessChain [[ARG1]]{{\[}}[[ZERO]]{{\]}} "iree.store_output"(%1, %arg1) : (i1, memref<i1>) -> () "iree.return"() : () -> ()
diff --git a/iree/compiler/Translation/SPIRV/test/gather.mlir b/iree/compiler/Translation/SPIRV/test/gather.mlir index d25551a..539f523 100644 --- a/iree/compiler/Translation/SPIRV/test/gather.mlir +++ b/iree/compiler/Translation/SPIRV/test/gather.mlir
@@ -15,18 +15,16 @@ // RUN: iree-opt -iree-index-computation -simplify-spirv-affine-exprs=false -convert-iree-to-spirv %s | IreeFileCheck %s module { - // CHECK: spv.globalVariable [[ARG0:@.*]] bind(0, 0) - // CHECK: spv.globalVariable [[ARG1:@.*]] bind(0, 1) - // CHECK: spv.globalVariable [[ARG2:@.*]] bind(0, 2) + // CHECK-LABEL: func @foo + // CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<5 x !spv.array<1 x !spv.array<10 x f32 [4]> [40]> [40]> [0]>, StorageBuffer> + // CHECK-SAME: [[ARG1:%.*]]: !spv.ptr<!spv.struct<i64 [0]>, StorageBuffer> func @foo(%arg0: memref<5x1x10xf32>, %arg1: memref<i64>, %arg2: memref<1x10xf32>) attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi64>, iree.executable.workload = dense<[10, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { - // CHECK: [[ADDRESS_ARG1:%.*]] = spv._address_of [[ARG1]] // CHECK: [[ZERO1:%.*]] = spv.constant 0 - // CHECK: [[LOAD_ADDRESS_ARG1:%.*]] = spv.AccessChain [[ADDRESS_ARG1]]{{\[}}[[ZERO1]]{{\]}} + // CHECK: [[LOAD_ADDRESS_ARG1:%.*]] = spv.AccessChain [[ARG1]]{{\[}}[[ZERO1]]{{\]}} // CHECK: [[INDEX1:%.*]] = spv.Load {{".*"}} [[LOAD_ADDRESS_ARG1]] - // CHECK: [[ADDRESS_ARG0:%.*]] = spv._address_of [[ARG0]] // CHECK: [[ZERO2:%.*]] = spv.constant 0 - // CHECK: {{%.*}} = spv.AccessChain [[ADDRESS_ARG0]]{{\[}}[[ZERO2]], [[INDEX1]] + // CHECK: {{%.*}} = spv.AccessChain [[ARG0]]{{\[}}[[ZERO2]], [[INDEX1]] %0 = iree.load_input(%arg0 : memref<5x1x10xf32>) : tensor<5x1x10xf32> %1 = iree.load_input(%arg1 : memref<i64>) : tensor<i64> %2 = "xla_hlo.gather"(%0, %1) {
diff --git a/iree/compiler/Translation/SPIRV/test/mulf.mlir b/iree/compiler/Translation/SPIRV/test/mulf.mlir index 318e7a1..e011328 100644 --- a/iree/compiler/Translation/SPIRV/test/mulf.mlir +++ b/iree/compiler/Translation/SPIRV/test/mulf.mlir
@@ -16,30 +16,28 @@ module { // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) - // CHECK-DAG: spv.globalVariable [[ARG1VAR:@.*]] bind(0, 1) - // CHECK-DAG: spv.globalVariable [[ARG2VAR:@.*]] bind(0, 2) + // CHECK: func @mul_1D + // CHECK-SAME: [[ARG0:%[a-zA-Z0-9]*]]: !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer> + // CHECK-SAME: [[ARG1:%[a-zA-Z0-9]*]]: !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer> + // CHECK-SAME: [[ARG2:%[a-zA-Z0-9]*]]: !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer> func @mul_1D(%arg0: memref<4xf32>, %arg1: memref<4xf32>, %arg2: memref<4xf32>) attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[GLOBALIDPTR:%.*]] = spv._address_of [[GLOBALIDVAR]] // CHECK: [[GLOBALID:%.*]] = spv.Load "Input" [[GLOBALIDPTR]] // CHECK: [[GLOBALIDX:%.*]] = spv.CompositeExtract [[GLOBALID]][0 : i32] // CHECK: spv.selection - // CHECK: [[ARG0PTR:%.*]] = spv._address_of [[ARG0VAR]] // CHECK: [[ZERO1:%.*]] = spv.constant 0 : i32 - // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0PTR]]{{\[}}[[ZERO1]], [[GLOBALIDX]]{{\]}} + // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]{{\[}}[[ZERO1]], [[GLOBALIDX]]{{\]}} // CHECK: [[VAL1:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]] %0 = iree.load_input(%arg0 : memref<4xf32>) : tensor<4xf32> - // CHECK: [[ARG1PTR:%.*]] = spv._address_of [[ARG1VAR]] // CHECK: [[ZERO2:%.*]] = spv.constant 0 : i32 - // CHECK: [[ARG1LOADPTR:%.*]] = spv.AccessChain [[ARG1PTR]]{{\[}}[[ZERO2]], [[GLOBALIDX]]{{\]}} + // CHECK: [[ARG1LOADPTR:%.*]] = spv.AccessChain [[ARG1]]{{\[}}[[ZERO2]], [[GLOBALIDX]]{{\]}} // CHECK: [[VAL2:%.*]] = spv.Load "StorageBuffer" [[ARG1LOADPTR]] %1 = iree.load_input(%arg1 : memref<4xf32>) : tensor<4xf32> // CHECK: [[RESULT:%.*]] = spv.FMul [[VAL1]], [[VAL2]] %2 = mulf %0, %1 : tensor<4xf32> - // CHECK: [[ARG2PTR:%.*]] = spv._address_of [[ARG2VAR]] // CHECK: [[ZERO3:%.*]] = spv.constant 0 : i32 - // CHECK: [[ARG2STOREPTR:%.*]] = spv.AccessChain [[ARG2PTR]]{{\[}}[[ZERO3]], [[GLOBALIDX]]{{\]}} + // CHECK: [[ARG2STOREPTR:%.*]] = spv.AccessChain [[ARG2]]{{\[}}[[ZERO3]], [[GLOBALIDX]]{{\]}} // CHECK: spv.Store "StorageBuffer" [[ARG2STOREPTR]], [[RESULT]] iree.store_output(%2 : tensor<4xf32>, %arg2 : memref<4xf32>) iree.return
diff --git a/iree/compiler/Translation/SPIRV/test/pad.mlir b/iree/compiler/Translation/SPIRV/test/pad.mlir index 945b645..3a848c5 100644 --- a/iree/compiler/Translation/SPIRV/test/pad.mlir +++ b/iree/compiler/Translation/SPIRV/test/pad.mlir
@@ -16,8 +16,9 @@ module { // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) - // CHECK-DAG: spv.globalVariable [[ARG1VAR:@.*]] bind(0, 1) + // CHECK: func @pad_zero_interior + // CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<4 x f32 [4]> [16]> [0]>, StorageBuffer> + // CHECK-SAME: [[ARG1:%.*]]: !spv.ptr<!spv.struct<!spv.array<18 x !spv.array<12 x f32 [4]> [48]> [0]>, StorageBuffer> func @pad_zero_interior(%arg0 : memref<12x4xf32>, %arg1 : memref<18x12xf32>) attributes {iree.executable.export, iree.executable.workload = dense<[12, 18, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[GLOBALIDPTR:%.*]] = spv._address_of [[GLOBALIDVAR]] @@ -25,13 +26,12 @@ // CHECK: [[GLOBALIDX:%.*]] = spv.CompositeExtract [[GLOBALID]][0 : i32] // CHECK: [[GLOBALIDY:%.*]] = spv.CompositeExtract [[GLOBALID]][1 : i32] // CHECK: spv.selection - // CHECK: [[ARG0PTR:%.*]] = spv._address_of [[ARG0VAR]] // CHECK: [[ZERO1:%.*]] = spv.constant 0 : i32 // CHECK: [[LOWER_PAD0:%.*]] = spv.constant -4 : i32 // CHECK: [[INDEX0:%.*]] = spv.IAdd [[GLOBALIDY]], [[LOWER_PAD0]] : i32 // CHECK: [[LOWER_PAD1:%.*]] = spv.constant -5 : i32 // CHECK: [[INDEX1:%.*]] = spv.IAdd [[GLOBALIDX]], [[LOWER_PAD1]] : i32 - // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0PTR]]{{\[}}[[ZERO1]], [[INDEX0]], [[INDEX1]]{{\]}} + // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]{{\[}}[[ZERO1]], [[INDEX0]], [[INDEX1]]{{\]}} // CHECK: [[INPUTVAL:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]] : f32 %0 = iree.load_input(%arg0 : memref<12x4xf32>) : tensor<12x4xf32> @@ -71,8 +71,8 @@ module { // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) - // CHECK-DAG: spv.globalVariable [[ARG1VAR:@.*]] bind(0, 1) + // CHECK: func @pad_no_op + // CHECK-SAME: [[ARG0:%[a-zA-Z0-9]*]]: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<4 x f32 [4]> [16]> [0]>, StorageBuffer> func @pad_no_op(%arg0 : memref<12x4xf32>, %arg1 : memref<12x4xf32>) attributes {iree.executable.export, iree.executable.workload = dense<[4, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[GLOBALIDPTR:%.*]] = spv._address_of [[GLOBALIDVAR]] @@ -80,9 +80,8 @@ // CHECK: [[GLOBALIDX:%.*]] = spv.CompositeExtract [[GLOBALID]][0 : i32] // CHECK: [[GLOBALIDY:%.*]] = spv.CompositeExtract [[GLOBALID]][1 : i32] // CHECK: spv.selection - // CHECK: [[ARG0PTR:%.*]] = spv._address_of [[ARG0VAR]] // CHECK: [[ZERO1:%.*]] = spv.constant 0 : i32 - // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0PTR]]{{\[}}[[ZERO1]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} + // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]{{\[}}[[ZERO1]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} // CHECK: [[INPUTVAL:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]] : f32 %0 = iree.load_input(%arg0 : memref<12x4xf32>) : tensor<12x4xf32> // CHECK: [[PADVAL:%.*]] = spv.constant 0.000000e+00 : f32 @@ -100,8 +99,8 @@ module { // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) - // CHECK-DAG: spv.globalVariable [[ARG1VAR:@.*]] bind(0, 1) + // CHECK: func @pad_zero_interior + // CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<4 x f32 [4]> [16]> [0]>, StorageBuffer> func @pad_zero_interior(%arg0 : memref<12x4xf32>, %arg1 : memref<29x18xf32>) attributes {iree.executable.export, iree.executable.workload = dense<[18, 29, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[GLOBALIDPTR:%.*]] = spv._address_of [[GLOBALIDVAR]] @@ -109,7 +108,6 @@ // CHECK: [[GLOBALIDX:%.*]] = spv.CompositeExtract [[GLOBALID]][0 : i32] // CHECK: [[GLOBALIDY:%.*]] = spv.CompositeExtract [[GLOBALID]][1 : i32] // CHECK: spv.selection - // CHECK: [[ARG0PTR:%.*]] = spv._address_of [[ARG0VAR]] // CHECK: [[ZERO1:%.*]] = spv.constant 0 : i32 // CHECK: [[INTERIOR0:%.*]] = spv.constant 2 : i32 // CHECK: [[SUB_PAD0:%.*]] = spv.SDiv [[GLOBALIDY]], [[INTERIOR0]] : i32 @@ -119,7 +117,7 @@ // CHECK: [[SUB_PAD1:%.*]] = spv.IAdd [[GLOBALIDX]], [[LOWER_PAD1]] : i32 // CHECK: [[INTERIOR1:%.*]] = spv.constant 3 : i32 // CHECK: [[INDEX1:%.*]] = spv.SDiv [[SUB_PAD1]], [[INTERIOR1]] : i32 - // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0PTR]]{{\[}}[[ZERO1]], [[INDEX0]], [[INDEX1]]{{\]}} + // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]{{\[}}[[ZERO1]], [[INDEX0]], [[INDEX1]]{{\]}} // CHECK: [[INPUTVAL:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]] : f32 %0 = iree.load_input(%arg0 : memref<12x4xf32>) : tensor<12x4xf32>
diff --git a/iree/compiler/Translation/SPIRV/test/reshape.mlir b/iree/compiler/Translation/SPIRV/test/reshape.mlir index f8f9a89..ac05794 100644 --- a/iree/compiler/Translation/SPIRV/test/reshape.mlir +++ b/iree/compiler/Translation/SPIRV/test/reshape.mlir
@@ -16,7 +16,8 @@ module { // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<24 x !spv.array<21 x i32 [4]> [84]> [0]>, StorageBuffer> + // CHECK: func @reshape_2D_2D + // CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<24 x !spv.array<21 x i32 [4]> [84]> [0]>, StorageBuffer> func @reshape_2D_2D(%arg0: memref<24x21xi32>, %arg1: memref<12x42xi32>) attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[GLOBALIDPTR:%.*]] = spv._address_of [[GLOBALIDVAR]] @@ -24,7 +25,6 @@ // CHECK: [[GLOBALIDX:%.*]] = spv.CompositeExtract [[GLOBALID]][0 : i32] // CHECK: [[GLOBALIDY:%.*]] = spv.CompositeExtract [[GLOBALID]][1 : i32] // CHECK: spv.selection - // CHECK-DAG: [[ARG0PTR:%.*]] = spv._address_of [[ARG0VAR]] // CHECK-DAG: [[ZERO1:%.*]] = spv.constant 0 : i32 // CHECK-DAG: [[STRIDE0:%.*]] = spv.constant 2 : i32 // CHECK-DAG: [[L1:%.*]] = spv.IMul [[GLOBALIDY]], [[STRIDE0]] : i32 @@ -32,7 +32,7 @@ // CHECK-DAG: [[L2:%.*]] = spv.SDiv [[GLOBALIDX]], [[STRIDE2]] : i32 // CHECK-DAG: [[INDEX0:%.*]] = spv.IAdd [[L1]], [[L2]] : i32 // CHECK-DAG: [[INDEX1:%.*]] = spv.SMod [[GLOBALIDX]], [[STRIDE2]] : i32 - // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0PTR]]{{\[}}[[ZERO1]], [[INDEX0]], [[INDEX1]]{{\]}} + // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]{{\[}}[[ZERO1]], [[INDEX0]], [[INDEX1]]{{\]}} %0 = iree.load_input(%arg0 : memref<24x21xi32>) : tensor<24x21xi32> %1 = "xla_hlo.reshape"(%0) : (tensor<24x21xi32>) -> tensor<12x42xi32> iree.store_output(%1 : tensor<12x42xi32>, %arg1 : memref<12x42xi32>) @@ -44,7 +44,8 @@ module { // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<4 x !spv.array<6 x !spv.array<21 x i32 [4]> [84]> [504]> [0]>, StorageBuffer> + // CHECK: func @reshape_3D_2D + // CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<4 x !spv.array<6 x !spv.array<21 x i32 [4]> [84]> [504]> [0]>, StorageBuffer> func @reshape_3D_2D(%arg0: memref<4x6x21xi32>, %arg1: memref<12x42xi32>) attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: spv.selection @@ -54,8 +55,7 @@ // CHECK-DAG: [[STRIDE2:%.*]] = spv.constant 21 : i32 // CHECK-DAG: [[I2:%.*]] = spv.SDiv [[I1]], [[STRIDE2]] : i32 // CHECK-DAG: [[I3:%.*]] = spv.SMod [[I1]], [[STRIDE2]] : i32 - // CHECK-DAG: [[ARG0PTR:%.*]] = spv._address_of [[ARG0VAR]] - // CHECK-DAG: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0PTR]]{{\[}}[[ZERO1]], [[I0]], [[I2]], [[I3]]{{\]}} + // CHECK-DAG: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]{{\[}}[[ZERO1]], [[I0]], [[I2]], [[I3]]{{\]}} %0 = iree.load_input(%arg0 : memref<4x6x21xi32>) : tensor<4x6x21xi32> %1 = "xla_hlo.reshape"(%0) : (tensor<4x6x21xi32>) -> tensor<12x42xi32> iree.store_output(%1 : tensor<12x42xi32>, %arg1 : memref<12x42xi32>) @@ -67,8 +67,9 @@ module { // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<24 x !spv.array<21 x i32 [4]> [84]> [0]>, StorageBuffer> - // CHECK-DAG: spv.globalVariable [[ARG1VAR:@.*]] bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<6 x !spv.array<7 x i32 [4]> [28]> [168]> [0]>, StorageBuffer> + // CHECK: func @reshape_2D_3D + // CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<24 x !spv.array<21 x i32 [4]> [84]> [0]>, StorageBuffer> + // CHECK-SAME: [[ARG1:%.*]]: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<6 x !spv.array<7 x i32 [4]> [28]> [168]> [0]>, StorageBuffer> func @reshape_2D_3D(%arg0: memref<24x21xi32>, %arg1: memref<12x6x7xi32>) attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[GLOBALIDPTR:%.*]] = spv._address_of [[GLOBALIDVAR]] @@ -76,7 +77,6 @@ // CHECK: [[GLOBALIDX:%.*]] = spv.CompositeExtract [[GLOBALID]][0 : i32] // CHECK: [[GLOBALIDY:%.*]] = spv.CompositeExtract [[GLOBALID]][1 : i32] // CHECK: spv.selection - // CHECK: [[ARG0PTR:%.*]] = spv._address_of [[ARG0VAR]] // CHECK-DAG: [[ZERO:%.*]] = spv.constant 0 : i32 // CHECK-DAG: [[FORTYTWO:%.*]] = spv.constant 42 : i32 // CHECK-DAG: [[I1:%.*]] = spv.SDiv [[GLOBALIDX]], [[FORTYTWO]] : i32 @@ -92,10 +92,9 @@ // CHECK: [[TWENTYONE:%.*]] = spv.constant 21 : i32 // CHECK-DAG: [[I11:%.*]] = spv.SDiv [[I10]], [[TWENTYONE]] : i32 // CHECK-DAG: [[I12:%.*]] = spv.SMod [[I10]], [[TWENTYONE]] : i32 - // CHECK: {{%.*}} = spv.AccessChain [[ARG0PTR]]{{\[}}[[ZERO]], [[I11]], [[I12]] - // CHECK-DAG: [[ARG1PTR:%.*]] = spv._address_of [[ARG1VAR]] + // CHECK: {{%.*}} = spv.AccessChain [[ARG0]]{{\[}}[[ZERO]], [[I11]], [[I12]] // CHECK-DAG: [[ZERO2:%.*]] = spv.constant 0 : i32 - // CHECK: {{%.*}} = spv.AccessChain [[ARG1PTR]]{{\[}}[[ZERO2]], [[I2]], [[I6]], [[I9]]{{\]}} + // CHECK: {{%.*}} = spv.AccessChain [[ARG1]]{{\[}}[[ZERO2]], [[I2]], [[I6]], [[I9]]{{\]}} %0 = iree.load_input(%arg0 : memref<24x21xi32>) : tensor<24x21xi32> %1 = "xla_hlo.reshape"(%0) : (tensor<24x21xi32>) -> tensor<12x6x7xi32> iree.store_output(%1 : tensor<12x6x7xi32>, %arg1 : memref<12x6x7xi32>)
diff --git a/iree/compiler/Translation/SPIRV/test/reshape_dropdims.mlir b/iree/compiler/Translation/SPIRV/test/reshape_dropdims.mlir index 7a9a9f2..8015216 100644 --- a/iree/compiler/Translation/SPIRV/test/reshape_dropdims.mlir +++ b/iree/compiler/Translation/SPIRV/test/reshape_dropdims.mlir
@@ -17,9 +17,8 @@ module { // CHECK:spv.module "Logical" "GLSL450" // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) - // CHECK-DAG: spv.globalVariable [[ARG1VAR:@.*]] bind(0, 1) - // CHECK: func [[FN:@reshape_4D_3D]]() + // CHECK: func [[FN:@reshape_4D_3D]] + // CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<42 x !spv.array<1 x i32 [4]> [4]> [168]> [0]>, StorageBuffer> func @reshape_4D_3D(%arg0: memref<12x42x1xi32>, %arg1: memref<12x42xi32>) attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[GLOBALIDPTR:%.*]] = spv._address_of [[GLOBALIDVAR]] @@ -27,10 +26,9 @@ // CHECK: [[GLOBALIDX:%.*]] = spv.CompositeExtract [[GLOBALID]][0 : i32] // CHECK: [[GLOBALIDY:%.*]] = spv.CompositeExtract [[GLOBALID]][1 : i32] // CHECK: spv.selection - // CHECK: [[ARG0PTR:%.*]] = spv._address_of [[ARG0VAR]] // CHECK: [[ZERO1:%.*]] = spv.constant 0 : i32 // CHECK: [[ZERO2:%.*]] = spv.constant 0 : i32 - // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0PTR]]{{\[}}[[ZERO1]], [[GLOBALIDY]], [[GLOBALIDX]], [[ZERO2]]{{\]}} + // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]{{\[}}[[ZERO1]], [[GLOBALIDY]], [[GLOBALIDX]], [[ZERO2]]{{\]}} %0 = iree.load_input(%arg0 : memref<12x42x1xi32>) : tensor<12x42x1xi32> %1 = "xla_hlo.reshape"(%0) : (tensor<12x42x1xi32>) -> tensor<12x42xi32> iree.store_output(%1 : tensor<12x42xi32>, %arg1 : memref<12x42xi32>) @@ -43,9 +41,8 @@ module { // CHECK:spv.module "Logical" "GLSL450" // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) - // CHECK-DAG: spv.globalVariable [[ARG1VAR:@.*]] bind(0, 1) - // CHECK: func [[FN:@reshape_4D_2D]]() + // CHECK: func [[FN:@reshape_4D_2D]] + // CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<42 x !spv.array<1 x !spv.array<1 x i32 [4]> [4]> [4]> [168]> [0]>, StorageBuffer> func @reshape_4D_2D(%arg0: memref<12x42x1x1xi32>, %arg1: memref<12x42xi32>) attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[GLOBALIDPTR:%.*]] = spv._address_of [[GLOBALIDVAR]] @@ -53,10 +50,9 @@ // CHECK: [[GLOBALIDX:%.*]] = spv.CompositeExtract [[GLOBALID]][0 : i32] // CHECK: [[GLOBALIDY:%.*]] = spv.CompositeExtract [[GLOBALID]][1 : i32] // CHECK: spv.selection - // CHECK: [[ARG0PTR:%.*]] = spv._address_of [[ARG0VAR]] // CHECK: [[ZERO1:%.*]] = spv.constant 0 : i32 // CHECK: [[ZERO2:%.*]] = spv.constant 0 : i32 - // CHECK: {{%.*}} = spv.AccessChain [[ARG0PTR]]{{\[}}[[ZERO1]], [[GLOBALIDY]], [[GLOBALIDX]], [[ZERO2]], [[ZERO2]]{{\]}} + // CHECK: {{%.*}} = spv.AccessChain [[ARG0]]{{\[}}[[ZERO1]], [[GLOBALIDY]], [[GLOBALIDX]], [[ZERO2]], [[ZERO2]]{{\]}} %0 = iree.load_input(%arg0 : memref<12x42x1x1xi32>) : tensor<12x42x1x1xi32> %1 = "xla_hlo.reshape"(%0) : (tensor<12x42x1x1xi32>) -> tensor<12x42xi32> iree.store_output(%1 : tensor<12x42xi32>, %arg1 : memref<12x42xi32>) @@ -69,9 +65,9 @@ module { // CHECK:spv.module "Logical" "GLSL450" // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) - // CHECK-DAG: spv.globalVariable [[ARG1VAR:@.*]] bind(0, 1) - // CHECK: func [[FN:@reshape_2D_4D]]() + // CHECK: func [[FN:@reshape_2D_4D]] + // CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<42 x i32 [4]> [168]> [0]>, StorageBuffer> + // CHECK-SAME: [[ARG1:%.*]]: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<42 x !spv.array<1 x !spv.array<1 x i32 [4]> [4]> [4]> [168]> [0]>, StorageBuffer> func @reshape_2D_4D(%arg0: memref<12x42xi32>, %arg1: memref<12x42x1x1xi32>) attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[GLOBALIDPTR:%.*]] = spv._address_of [[GLOBALIDVAR]] @@ -79,15 +75,13 @@ // CHECK: [[GLOBALIDX:%.*]] = spv.CompositeExtract [[GLOBALID]][0 : i32] // CHECK: [[GLOBALIDY:%.*]] = spv.CompositeExtract [[GLOBALID]][1 : i32] // CHECK: spv.selection - // CHECK: [[ARG0PTR:%.*]] = spv._address_of [[ARG0VAR]] // CHECK: [[ZERO1:%.*]] = spv.constant 0 : i32 - // CHECK: {{%.*}} = spv.AccessChain [[ARG0PTR]]{{\[}}[[ZERO1]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} + // CHECK: {{%.*}} = spv.AccessChain [[ARG0]]{{\[}}[[ZERO1]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} %0 = iree.load_input(%arg0 : memref<12x42xi32>) : tensor<12x42xi32> %1 = "xla_hlo.reshape"(%0) : (tensor<12x42xi32>) -> tensor<12x42x1x1xi32> - // CHECK: [[ARG1PTR:%.*]] = spv._address_of [[ARG1VAR]] // CHECK: [[ZERO2:%.*]] = spv.constant 0 : i32 // CHECK: [[ZERO3:%.*]] = spv.constant 0 : i32 - // CHECK: {{%.*}} = spv.AccessChain [[ARG1PTR]]{{\[}}[[ZERO2]], [[GLOBALIDY]], [[GLOBALIDX]], [[ZERO3]], [[ZERO3]]{{\]}} + // CHECK: {{%.*}} = spv.AccessChain [[ARG1]]{{\[}}[[ZERO2]], [[GLOBALIDY]], [[GLOBALIDX]], [[ZERO3]], [[ZERO3]]{{\]}} iree.store_output(%1 : tensor<12x42x1x1xi32>, %arg1 : memref<12x42x1x1xi32>) iree.return } @@ -98,9 +92,9 @@ module { // CHECK:spv.module "Logical" "GLSL450" // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) - // CHECK-DAG: spv.globalVariable [[ARG1VAR:@.*]] bind(0, 1) - // CHECK: func [[FN:@reshape_2D_4D]]() + // CHECK: func [[FN:@reshape_2D_4D]] + // CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<42 x i32 [4]> [168]> [0]>, StorageBuffer> + // CHECK-SAME: [[ARG1:%.*]]: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<1 x !spv.array<1 x !spv.array<42 x i32 [4]> [168]> [168]> [168]> [0]>, StorageBuffer> func @reshape_2D_4D(%arg0: memref<12x42xi32>, %arg1: memref<12x1x1x42xi32>) attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { %0 = iree.load_input(%arg0 : memref<12x42xi32>) : tensor<12x42xi32> @@ -109,14 +103,12 @@ // CHECK: [[GLOBALIDX:%.*]] = spv.CompositeExtract [[GLOBALID]][0 : i32] // CHECK: [[GLOBALIDY:%.*]] = spv.CompositeExtract [[GLOBALID]][1 : i32] // CHECK: spv.selection - // CHECK: [[ARG0PTR:%.*]] = spv._address_of [[ARG0VAR]] // CHECK: [[ZERO1:%.*]] = spv.constant 0 : i32 - // CHECK: {{%.*}} = spv.AccessChain [[ARG0PTR]]{{\[}}[[ZERO1]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} + // CHECK: {{%.*}} = spv.AccessChain [[ARG0]]{{\[}}[[ZERO1]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} %1 = "xla_hlo.reshape"(%0) : (tensor<12x42xi32>) -> tensor<12x1x1x42xi32> - // CHECK: [[ARG1PTR:%.*]] = spv._address_of [[ARG1VAR]] // CHECK: [[ZERO2:%.*]] = spv.constant 0 : i32 // CHECK: [[ZERO3:%.*]] = spv.constant 0 : i32 - // CHECK: {{%.*}} = spv.AccessChain [[ARG1PTR]]{{\[}}[[ZERO2]], [[GLOBALIDY]], [[ZERO3]], [[ZERO3]], [[GLOBALIDX]]{{\]}} + // CHECK: {{%.*}} = spv.AccessChain [[ARG1]]{{\[}}[[ZERO2]], [[GLOBALIDY]], [[ZERO3]], [[ZERO3]], [[GLOBALIDX]]{{\]}} iree.store_output(%1 : tensor<12x1x1x42xi32>, %arg1 : memref<12x1x1x42xi32>) iree.return } @@ -127,9 +119,9 @@ module { // CHECK:spv.module "Logical" "GLSL450" // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) - // CHECK-DAG: spv.globalVariable [[ARG1VAR:@.*]] bind(0, 1) - // CHECK: func [[FN:@reshape_2D_4D]]() + // CHECK: func [[FN:@reshape_2D_4D]] + // CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<1 x !spv.array<1 x !spv.array<42 x i32 [4]> [168]> [168]> [168]> [0]>, StorageBuffer> + // CHECK-SAME: [[ARG1:%.*]]: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<42 x i32 [4]> [168]> [0]>, StorageBuffer> func @reshape_2D_4D(%arg0: memref<12x1x1x42xi32>, %arg1: memref<12x42xi32>) attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[GLOBALIDPTR:%.*]] = spv._address_of [[GLOBALIDVAR]] @@ -137,15 +129,13 @@ // CHECK: [[GLOBALIDX:%.*]] = spv.CompositeExtract [[GLOBALID]][0 : i32] // CHECK: [[GLOBALIDY:%.*]] = spv.CompositeExtract [[GLOBALID]][1 : i32] // CHECK: spv.selection - // CHECK: [[ARG0PTR:%.*]] = spv._address_of [[ARG0VAR]] // CHECK: [[ZERO1:%.*]] = spv.constant 0 : i32 // CHECK: [[ZERO2:%.*]] = spv.constant 0 : i32 - // CHECK: {{%.*}} = spv.AccessChain [[ARG0PTR]]{{\[}}[[ZERO1]], [[GLOBALIDY]], [[ZERO2]], [[ZERO2]], [[GLOBALIDX]]{{\]}} + // CHECK: {{%.*}} = spv.AccessChain [[ARG0]]{{\[}}[[ZERO1]], [[GLOBALIDY]], [[ZERO2]], [[ZERO2]], [[GLOBALIDX]]{{\]}} %0 = iree.load_input(%arg0 : memref<12x1x1x42xi32>) : tensor<12x1x1x42xi32> %1 = "xla_hlo.reshape"(%0) : (tensor<12x1x1x42xi32>) -> tensor<12x42xi32> - // CHECK: [[ARG1PTR:%.*]] = spv._address_of [[ARG1VAR]] // CHECK: [[ZERO3:%.*]] = spv.constant 0 : i32 - // CHECK: {{%.*}} = spv.AccessChain [[ARG1PTR]]{{\[}}[[ZERO3]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} + // CHECK: {{%.*}} = spv.AccessChain [[ARG1]]{{\[}}[[ZERO3]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} iree.store_output(%1 : tensor<12x42xi32>, %arg1 : memref<12x42xi32>) iree.return }
diff --git a/iree/compiler/Translation/SPIRV/test/reverse.mlir b/iree/compiler/Translation/SPIRV/test/reverse.mlir index 9fda215..4a3f4f5 100644 --- a/iree/compiler/Translation/SPIRV/test/reverse.mlir +++ b/iree/compiler/Translation/SPIRV/test/reverse.mlir
@@ -16,8 +16,8 @@ module { // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) - // CHECK-DAG: spv.globalVariable [[ARG1VAR:@.*]] bind(0, 1) + // CHECK: func @reverse_2d + // CHECK-SAME: [[ARG0:%[a-zA-Z0-9]*]]: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<12 x f32 [4]> [48]> [0]>, StorageBuffer> func @reverse_2d(%arg0: memref<12x12xf32>, %arg1 : memref<12x12xf32>) attributes {iree.executable.export, iree.executable.workload = dense<[12, 12]> : tensor<2xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[GLOBALIDPTR:%.*]] = spv._address_of [[GLOBALIDVAR]] @@ -25,7 +25,6 @@ // CHECK: [[GLOBALIDX:%.*]] = spv.CompositeExtract [[GLOBALID]][0 : i32] // CHECK: [[GLOBALIDY:%.*]] = spv.CompositeExtract [[GLOBALID]][1 : i32] // CHECK: spv.selection - // CHECK: [[ARG0PTR:%.*]] = spv._address_of [[ARG0VAR]] // CHECK: [[ZERO:%.*]] = spv.constant 0 : i32 // CHECK: [[NEGATIVE_ONE:%.*]] = spv.constant -1 : i32 // CHECK: [[NEGATIVE_IDY:%.*]] = spv.IMul [[GLOBALIDY]], [[NEGATIVE_ONE]] : i32 @@ -33,7 +32,7 @@ // CHECK: [[REVERSE_IDY:%.*]] = spv.IAdd [[NEGATIVE_IDY]], [[ELEVEN]] : i32 // CHECK: [[NEGATIVE_IDX:%.*]] = spv.IMul [[GLOBALIDX]], [[NEGATIVE_ONE]] : i32 // CHECK: [[REVERSE_IDX:%.*]] = spv.IAdd [[NEGATIVE_IDX]], [[ELEVEN]] : i32 - // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0PTR]]{{\[}}[[ZERO]], [[REVERSE_IDY]], [[REVERSE_IDX]]{{\]}} + // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]{{\[}}[[ZERO]], [[REVERSE_IDY]], [[REVERSE_IDX]]{{\]}} // CHECK: [[VAL0:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]] : f32 %0 = iree.load_input(%arg0 : memref<12x12xf32>) : tensor<12x12xf32> %1 = "xla_hlo.reverse"(%0) {dimensions = dense<[1, 0]> : tensor<2xi64>} : (tensor<12x12xf32>) -> tensor<12x12xf32> @@ -46,8 +45,8 @@ module { // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) - // CHECK-DAG: spv.globalVariable [[ARG1VAR:@.*]] bind(0, 1) + // CHECK: func @reverse_3d + // CHECK-SAME: [[ARG0:%[a-zA-Z0-9]*]]: !spv.ptr<!spv.struct<!spv.array<3 x !spv.array<3 x !spv.array<3 x f32 [4]> [12]> [36]> [0]>, StorageBuffer> func @reverse_3d(%arg0: memref<3x3x3xf32>, %arg1 : memref<3x3x3xf32>) attributes {iree.executable.export, iree.executable.workload = dense<[3, 3, 3]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[GLOBALIDPTR:%.*]] = spv._address_of [[GLOBALIDVAR]] @@ -56,13 +55,12 @@ // CHECK: [[GLOBALIDY:%.*]] = spv.CompositeExtract [[GLOBALID]][1 : i32] // CHECK: [[GLOBALIDZ:%.*]] = spv.CompositeExtract [[GLOBALID]][2 : i32] // CHECK: spv.selection - // CHECK: [[ARG0PTR:%.*]] = spv._address_of [[ARG0VAR]] // CHECK: [[ZERO:%.*]] = spv.constant 0 : i32 // CHECK: [[NEGATIVE_ONE:%.*]] = spv.constant -1 : i32 // CHECK: [[NEGATIVE_IDY:%.*]] = spv.IMul [[GLOBALIDY]], [[NEGATIVE_ONE]] : i32 // CHECK: [[TWO:%.*]] = spv.constant 2 : i32 // CHECK: [[REVERSE_IDY:%.*]] = spv.IAdd [[NEGATIVE_IDY]], [[TWO]] : i32 - // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0PTR]]{{\[}}[[ZERO]], [[GLOBALIDZ]], [[REVERSE_IDY]], [[GLOBALIDX]]{{\]}} + // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]{{\[}}[[ZERO]], [[GLOBALIDZ]], [[REVERSE_IDY]], [[GLOBALIDX]]{{\]}} // CHECK: [[VAL0:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]] : f32 %0 = iree.load_input(%arg0 : memref<3x3x3xf32>) : tensor<3x3x3xf32> %1 = "xla_hlo.reverse"(%0) {dimensions = dense<[1]> : tensor<1xi64>} : (tensor<3x3x3xf32>) -> tensor<3x3x3xf32>
diff --git a/iree/compiler/Translation/SPIRV/test/slice.mlir b/iree/compiler/Translation/SPIRV/test/slice.mlir index a933d51..2a912a2 100644 --- a/iree/compiler/Translation/SPIRV/test/slice.mlir +++ b/iree/compiler/Translation/SPIRV/test/slice.mlir
@@ -20,8 +20,9 @@ // VAR0 = (x / 3) // VAR2 = (X mod 3) // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) - // CHECK-DAG: spv.globalVariable [[ARG1VAR:@.*]] bind(0, 1) + // CHECK: func @slice_unit_stride + // CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<6 x !spv.array<6 x f32 [4]> [24]> [0]>, StorageBuffer> + // CHECK-SAME: [[ARG1:%.*]]: !spv.ptr<!spv.struct<!spv.array<2 x !spv.array<3 x f32 [4]> [12]> [0]>, StorageBuffer> func @slice_unit_stride(%arg0: memref<6x6xf32>, %arg1: memref<2x3xf32>) attributes {iree.executable.export, iree.executable.workload = dense<[6, 1]> : tensor<2xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[GLOBALIDPTR:%.*]] = spv._address_of [[GLOBALIDVAR]] @@ -29,7 +30,6 @@ // CHECK: [[GLOBALIDX:%.*]] = spv.CompositeExtract [[GLOBALID]][0 : i32] // CHECK: [[GLOBALIDY:%.*]] = spv.CompositeExtract [[GLOBALID]][1 : i32] // CHECK: spv.selection - // CHECK: [[ARG0PTR:%.*]] = spv._address_of [[ARG0VAR]] // CHECK: [[ZERO:%.*]] = spv.constant 0 : i32 // CHECK: [[THREE:%.*]] = spv.constant 3 : i32 // CHECK: [[VAR0:%.*]] = spv.SDiv [[GLOBALIDX]], [[THREE]] : i32 @@ -38,11 +38,10 @@ // CHECK: [[VAR2:%.*]] = spv.SMod [[GLOBALIDX]], [[THREE]] : i32 // CHECK: [[ONE:%.*]] = spv.constant 1 : i32 // CHECK: [[VAR3:%.*]] = spv.IAdd [[VAR2]], [[ONE]] : i32 - // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0PTR]]{{\[}}[[ZERO]], [[VAR1]], [[VAR3]]{{\]}} + // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]{{\[}}[[ZERO]], [[VAR1]], [[VAR3]]{{\]}} // CHECK: [[VAL0:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]] : f32 - // CHECK: [[ARG1PTR:%.*]] = spv._address_of [[ARG1VAR]] // CHECK: [[ZERO2:%.*]] = spv.constant 0 : i32 - // CHECK: [[ARG1STOREPTR:%.*]] = spv.AccessChain [[ARG1PTR]]{{\[}}[[ZERO2]], [[VAR0]], [[VAR2]]{{\]}} + // CHECK: [[ARG1STOREPTR:%.*]] = spv.AccessChain [[ARG1]]{{\[}}[[ZERO2]], [[VAR0]], [[VAR2]]{{\]}} // CHECK: spv.Store "StorageBuffer" [[ARG1STOREPTR]], [[VAL0]] : f32 %0 = iree.load_input(%arg0 : memref<6x6xf32>) : tensor<6x6xf32> %1 = "xla_hlo.slice"(%0) {start_indices = dense<[2, 1]> : tensor<2xi64>, limit_indices = dense<[4, 4]> : tensor<2xi64>, strides = dense<[1, 1]> : tensor<2xi64>} : (tensor<6x6xf32>) -> tensor<2x3xf32> @@ -59,8 +58,9 @@ // VAR0 = (x / 3) // VAR2 = (x mod 3) // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) - // CHECK-DAG: spv.globalVariable [[ARG1VAR:@.*]] bind(0, 1) + // CHECK: func @slice_non_unit_stride + // CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<6 x !spv.array<6 x f32 [4]> [24]> [0]>, StorageBuffer> + // CHECK-SAME: [[ARG1:%.*]]: !spv.ptr<!spv.struct<!spv.array<2 x !spv.array<3 x f32 [4]> [12]> [0]>, StorageBuffer> func @slice_non_unit_stride(%arg0: memref<6x6xf32>, %arg1: memref<2x3xf32>) attributes {iree.executable.export, iree.executable.workload = dense<[6, 1]> : tensor<2xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[GLOBALIDPTR:%.*]] = spv._address_of [[GLOBALIDVAR]] @@ -68,7 +68,6 @@ // CHECK: [[GLOBALIDX:%.*]] = spv.CompositeExtract %1[0 : i32] : vector<3xi32> // CHECK: [[GLOBALIDY:%.*]] = spv.CompositeExtract %1[1 : i32] : vector<3xi32> // CHECK: spv.selection - // CHECK: [[ARG0PTR:%.*]] = spv._address_of [[ARG0VAR]] // CHECK: [[ZERO:%.*]] = spv.constant 0 : i32 // CHECK: [[THREE:%.*]] = spv.constant 3 : i32 // CHECK: [[VAR0:%.*]] = spv.SDiv [[GLOBALIDX]], [[THREE]] : i32 @@ -78,11 +77,10 @@ // CHECK: [[VAR3:%.*]] = spv.IMul [[VAR2]], [[TWO]] : i32 // CHECK: [[ONE:%.*]] = spv.constant 1 : i32 // CHECK: [[VAR4:%.*]] = spv.IAdd [[VAR3]], [[ONE]] : i32 - // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0PTR]]{{\[}}[[ZERO]], [[VAR1]], [[VAR4]]{{\]}} + // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]{{\[}}[[ZERO]], [[VAR1]], [[VAR4]]{{\]}} // CHECK: [[VAL0:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]] : f32 - // CHECK: [[ARG1PTR:%.*]] = spv._address_of [[ARG1VAR]] // CHECK: [[ZERO2:%.*]] = spv.constant 0 : i32 - // CHECK: [[ARG1STOREPTR:%.*]] = spv.AccessChain [[ARG1PTR]]{{\[}}[[ZERO2]], [[VAR0]], [[VAR2]]{{\]}} + // CHECK: [[ARG1STOREPTR:%.*]] = spv.AccessChain [[ARG1]]{{\[}}[[ZERO2]], [[VAR0]], [[VAR2]]{{\]}} // CHECK: spv.Store "StorageBuffer" [[ARG1STOREPTR]], [[VAL0]] : f32 %0 = iree.load_input(%arg0 : memref<6x6xf32>) : tensor<6x6xf32> %1 = "xla_hlo.slice"(%0) {start_indices = dense<[2, 1]> : tensor<2xi64>, limit_indices = dense<[4, 6]> : tensor<2xi64>, strides = dense<[1, 2]> : tensor<2xi64>} : (tensor<6x6xf32>) -> tensor<2x3xf32>
diff --git a/iree/compiler/Translation/SPIRV/test/transpose_add.mlir b/iree/compiler/Translation/SPIRV/test/transpose_add.mlir index 9b8ce26..3c2e230 100644 --- a/iree/compiler/Translation/SPIRV/test/transpose_add.mlir +++ b/iree/compiler/Translation/SPIRV/test/transpose_add.mlir
@@ -16,22 +16,20 @@ module { // CHECK-DAG: spv.globalVariable [[GLOBALIDVAR:@.*]] built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input> - // CHECK-DAG: spv.globalVariable [[ARG0VAR:@.*]] bind(0, 0) - // CHECK-DAG: spv.globalVariable [[ARG1VAR:@.*]] bind(0, 1) - func @transpose_add(%arg0: memref<12x12xf32>, %arg1: memref<12x12xf32>) + // CHECK: func @transpose_add + // CHECK-SAME: [[ARG0:%[a-zA-Z0-9]*]]: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<12 x f32 [4]> [48]> [0]>, StorageBuffer> + func @transpose_add(%arg0: memref<12x12xf32>, %arg1: memref<12x12xf32>) attributes {iree.executable.export, iree.executable.workload = dense<[12, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { // CHECK: [[GLOBALIDPTR:%.*]] = spv._address_of [[GLOBALIDVAR]] // CHECK: [[GLOBALID:%.*]] = spv.Load "Input" [[GLOBALIDPTR]] // CHECK: [[GLOBALIDX:%.*]] = spv.CompositeExtract [[GLOBALID]][0 : i32] // CHECK: [[GLOBALIDY:%.*]] = spv.CompositeExtract [[GLOBALID]][1 : i32] // CHECK: spv.selection - // CHECK: [[ARG0PTR:%.*]] = spv._address_of [[ARG0VAR]] // CHECK: [[ZERO1:%.*]] = spv.constant 0 : i32 - // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0PTR]]{{\[}}[[ZERO1]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} + // CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]{{\[}}[[ZERO1]], [[GLOBALIDY]], [[GLOBALIDX]]{{\]}} // CHECK: [[VAL1:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]] - // CHECK: [[ARG1PTR:%.*]] = spv._address_of [[ARG0VAR]] // CHECK: [[ZERO2:%.*]] = spv.constant 0 : i32 - // CHECK: [[ARG1LOADPTR:%.*]] = spv.AccessChain [[ARG1PTR]]{{\[}}[[ZERO2]], [[GLOBALIDX]], [[GLOBALIDY]]{{\]}} + // CHECK: [[ARG1LOADPTR:%.*]] = spv.AccessChain [[ARG0]]{{\[}}[[ZERO2]], [[GLOBALIDX]], [[GLOBALIDY]]{{\]}} // CHECK: [[VAL2:%.*]] = spv.Load "StorageBuffer" [[ARG1LOADPTR]] %0 = iree.load_input(%arg0 : memref<12x12xf32>) : tensor<12x12xf32> %1 = "xla_hlo.transpose"(%0) {permutation = dense<[1, 0]> : tensor<2xi64>} : (tensor<12x12xf32>) -> tensor<12x12xf32>