[mlir] Flip Async/GPU/OpenACC/OpenMP to use Both accessors

[rriddle]

Summary

This allows for incrementally updating the old API usages without
needing to update everything at once. These will be left on Both
for a little bit and then flipped to prefixed when all APIs have been
updated.

flang/lib/Lower/OpenACC.cpp

createLoopOp(Fortran::lower::AbstractConverter &converter,
   addOperand(operands, operandSegments, vectorNum);
   addOperands(operands, operandSegments, tileOperands);
   addOperands(operands, operandSegments, privateOperands);
   addOperands(operands, operandSegments, reductionOperands);
 
   auto loopOp = createRegionOp<mlir::acc::LoopOp, mlir::acc::YieldOp>(
       firOpBuilder, currentLocation, operands, operandSegments);
 
-  loopOp->setAttr(mlir::acc::LoopOp::getExecutionMappingAttrName(),
-                  firOpBuilder.getI64IntegerAttr(executionMapping));
+  loopOp.setExecMappingAttr(firOpBuilder.getI64IntegerAttr(executionMapping));
 
   // Lower clauses mapped to attributes
   for (const Fortran::parser::AccClause &clause : accClauseList.v) {
     if (const auto *collapseClause =
             std::get_if<Fortran::parser::AccClause::Collapse>(&clause.u)) {
       const auto *expr = Fortran::semantics::GetExpr(collapseClause->v);
       const std::optional<int64_t> collapseValue =
           Fortran::evaluate::ToInt64(*expr);
       if (collapseValue) {
-        loopOp->setAttr(mlir::acc::LoopOp::getCollapseAttrName(),
-                        firOpBuilder.getI64IntegerAttr(*collapseValue));
+        loopOp.setCollapseAttr(firOpBuilder.getI64IntegerAttr(*collapseValue));
       }
     } else if (std::get_if<Fortran::parser::AccClause::Seq>(&clause.u)) {
-      loopOp->setAttr(mlir::acc::LoopOp::getSeqAttrName(),
-                      firOpBuilder.getUnitAttr());
+      loopOp.setSeqAttr(firOpBuilder.getUnitAttr());
     } else if (std::get_if<Fortran::parser::AccClause::Independent>(
                    &clause.u)) {
-      loopOp->setAttr(mlir::acc::LoopOp::getIndependentAttrName(),
-                      firOpBuilder.getUnitAttr());
+      loopOp.setIndependentAttr(firOpBuilder.getUnitAttr());
     } else if (std::get_if<Fortran::parser::AccClause::Auto>(&clause.u)) {
-      loopOp->setAttr(mlir::acc::LoopOp::getAutoAttrName(),
+      loopOp->setAttr(mlir::acc::LoopOp::getAutoAttrStrName(),
                       firOpBuilder.getUnitAttr());
     }
   }
   return loopOp;
 }
 
 static void genACC(Fortran::lower::AbstractConverter &converter,
                    Fortran::semantics::SemanticsContext &semanticsContext,

mlir/include/mlir/Dialect/Async/IR/AsyncDialect.td

let extraClassDeclaration = [{
     /// blocking async.runtime.await ops. Only useful in combination with
     /// 'eliminate-blocking-await-ops' option, which in absence of this
     /// attribute might convert a func to a coroutine.
     static constexpr StringRef kAllowedToBlockAttrName =
         "async.allowed_to_block";
   }];
 
   // TODO: Prefixed form overlaps with generated names, update before flipping.
-  let emitAccessorPrefix = kEmitAccessorPrefix_Raw;
+  let emitAccessorPrefix = kEmitAccessorPrefix_Both;
 }
 
 #endif // ASYNC_DIALECT_TD

mlir/include/mlir/Dialect/Async/IR/AsyncOps.td

`async.execute` operation takes `async.token` dependencies and `async.value`
     ```
 
     In the example above asynchronous execution starts only after dependency
     token and value argument become ready. Unwrapped value passed to the
     attached body region as an %unwrapped value of f32 type.
   }];
 
   let arguments = (ins Variadic<Async_TokenType>:$dependencies,
-                       Variadic<Async_AnyValueOrTokenType>:$operands);
+                       Variadic<Async_AnyValueOrTokenType>:$bodyOperands);
 
   let results = (outs Async_TokenType:$token,
-                      Variadic<Async_ValueType>:$results);
-  let regions = (region SizedRegion<1>:$body);
+                      Variadic<Async_ValueType>:$bodyResults);
+  let regions = (region SizedRegion<1>:$bodyRegion);
 
   let hasCustomAssemblyFormat = 1;
   let skipDefaultBuilders = 1;
   let hasRegionVerifier = 1;
   let builders = [
     OpBuilder<(ins "TypeRange":$resultTypes, "ValueRange":$dependencies,
       "ValueRange":$operands,
       CArg<"function_ref<void(OpBuilder &, Location, ValueRange)>",

mlir/include/mlir/Dialect/GPU/IR/GPUBase.td

let extraClassDeclaration = [{
     /// space.
     static unsigned getPrivateAddressSpace() { return 5; }
   }];
 
   let dependentDialects = ["arith::ArithmeticDialect"];
   let useDefaultAttributePrinterParser = 1;
   let useDefaultTypePrinterParser = 1;
 
-  // TODO: This has overlapping accessors with generated when switched to
-  // prefixed. Fix and update to _Both & then _Prefixed.
-  let emitAccessorPrefix = kEmitAccessorPrefix_Raw;
+  // TODO: Update this to _Prefixed.
+  let emitAccessorPrefix = kEmitAccessorPrefix_Both;
 }
 
 def GPU_AsyncToken : DialectType<
   GPU_Dialect, CPred<"$_self.isa<::mlir::gpu::AsyncTokenType>()">, "async token type">,
              BuildableType<"mlir::gpu::AsyncTokenType::get($_builder.getContext())">;
 
 // Predicat to check if type is gpu::MMAMatrixType.
 def IsMMAMatrixTypePred : CPred<"$_self.isa<::mlir::gpu::MMAMatrixType>()">;
InterfaceMethod<[{
       "void", "addAsyncDependency", (ins "Value":$token),
       [{}], [{
         ::mlir::gpu::addAsyncDependency(this->getOperation(), token);
       }]
     >,
     InterfaceMethod<[{
         Query the result that represents the async token to depend on.
       }],
-      "OpResult", "getAsyncToken", (ins), [{}], [{
-        ConcreteOp op = cast<ConcreteOp>(this->getOperation());
-        return op.asyncToken().template dyn_cast_or_null<OpResult>();
-      }]
+      "Value", "getAsyncToken"
     >
   ];
 }
 
 //===----------------------------------------------------------------------===//
 // GPU Attributes.
 //===----------------------------------------------------------------------===//
 
 class GPU_Attr<string attrName, string attrMnemonic, list<Trait> traits = []>
     : AttrDef<GPU_Dialect, attrName, traits> {
   let mnemonic = attrMnemonic;
 }
 
 #endif // GPU_BASE

mlir/include/mlir/Dialect/GPU/IR/GPUOps.td

let extraClassDeclaration = [{
     BlockArgument addPrivateAttribution(Type type, Location loc);
 
     /// Returns the name of the attribute containing the number of buffers
     /// located in the workgroup memory.
     static StringRef getNumWorkgroupAttributionsAttrName() {
       return "workgroup_attributions";
     }
 
-    /// Returns the type of this function.
-    /// FIXME: Remove when GPU uses prefixed accessors.
-    FunctionType getFunctionType() { return function_type(); }
-
     /// Returns the argument types of this function.
     ArrayRef<Type> getArgumentTypes() { return getFunctionType().getInputs(); }
 
     /// Returns the result types of this function.
     ArrayRef<Type> getResultTypes() { return getFunctionType().getResults(); }
 
     /// Returns the keywords used in the custom syntax for this Op.
     static StringRef getWorkgroupKeyword() { return "workgroup"; }

 
 def GPU_LaunchFuncOp : GPU_Op<"launch_func",
                               [GPU_AsyncOpInterface, AttrSizedOperandSegments]>,
     Arguments<(ins Variadic<GPU_AsyncToken>:$asyncDependencies,
                SymbolRefAttr:$kernel,
                Index:$gridSizeX, Index:$gridSizeY, Index:$gridSizeZ,
                Index:$blockSizeX, Index:$blockSizeY, Index:$blockSizeZ,
                Optional<I32>:$dynamicSharedMemorySize,
-               Variadic<AnyType>:$operands)>,
+               Variadic<AnyType>:$kernelOperands)>,
     Results<(outs Optional<GPU_AsyncToken>:$asyncToken)> {
   let summary = "Launches a function as a GPU kernel";
 
   let description = [{
     Launch a kernel function on the specified grid of thread blocks.
     `gpu.launch` operations are lowered to `gpu.launch_func` operations by
     outlining the kernel body into a function in a dedicated module, which
     reflects the separate compilation process. The kernel function is required
let extraClassDeclaration = [{
 
     /// Get the SSA values passed as operands to specify the block size.
     KernelDim3 getBlockSizeOperandValues();
 
     // This needs to quietly verify if attributes with names defined below are
     // present since it is run before the verifier of this op.
     friend LogicalResult GPUDialect::verifyOperationAttribute(Operation *,
                                                               NamedAttribute);
-
-    /// The name of the symbol reference attribute specifying the kernel to launch.
-    static StringRef getKernelAttrName() { return "kernel"; }
   }];
 
   let assemblyFormat = [{
       custom<AsyncDependencies>(type($asyncToken), $asyncDependencies)
       $kernel
       `blocks` `in` ` ` `(`$gridSizeX`,` $gridSizeY`,` $gridSizeZ`)`
       `threads` `in` ` ` `(`$blockSizeX`,` $blockSizeY`,` $blockSizeZ`)`
       (`dynamic_shared_memory_size` $dynamicSharedMemorySize^)?
-      custom<LaunchFuncOperands>($operands, type($operands)) attr-dict
+      custom<LaunchFuncOperands>($kernelOperands, type($kernelOperands)) attr-dict
   }];
   let hasVerifier = 1;
 }
 
 def GPU_LaunchOp : GPU_Op<"launch", [
       AutomaticAllocationScope, AttrSizedOperandSegments, GPU_AsyncOpInterface,
       DeclareOpInterfaceMethods<InferIntRangeInterface>]>,
     Arguments<(ins Variadic<GPU_AsyncToken>:$asyncDependencies,

 }
 def GPU_ShuffleModeAttr : EnumAttr<GPU_Dialect, GPU_ShuffleMode,
                                    "shuffle_mode">;
 
 def I32OrF32 : TypeConstraint<Or<[I32.predicate, F32.predicate]>,
                                  "i32 or f32">;
 
 def GPU_ShuffleOp : GPU_Op<
-    "shuffle", [NoSideEffect, AllTypesMatch<["value", "result"]>]>,
+    "shuffle", [NoSideEffect, AllTypesMatch<["value", "shuffleResult"]>]>,
     Arguments<(ins I32OrF32:$value, I32:$offset, I32:$width,
                GPU_ShuffleModeAttr:$mode)>,
-    Results<(outs I32OrF32:$result, I1:$valid)> {
+    Results<(outs I32OrF32:$shuffleResult, I1:$valid)> {
   let summary = "Shuffles values within a subgroup.";
   let description = [{
     The "shuffle" op moves values to a different invocation within the same
     subgroup.
 
     Example:
 
     ```mlir
```
       gpu.func {}
         ...
       gpu.module_end
     }
 
     ```
   }];
   let builders = [OpBuilder<(ins "StringRef":$name)>];
-  let regions = (region SizedRegion<1>:$body);
+  let regions = (region SizedRegion<1>:$bodyRegion);
   let hasCustomAssemblyFormat = 1;
 
   // We need to ensure the block inside the region is properly terminated;
   // the auto-generated builders do not guarantee that.
   let skipDefaultBuilders = 1;
 }
 
 def GPU_ModuleEndOp : GPU_Op<"module_end", [
let description = [{
     Since the operation is elementwise and the matrix type must match, the
     matrix elements are processed independently of the matrix layout.
 
     This op is meant to be used along with `gpu.subgroup_mma_compute`.
 
     Example:
 
     ```mlir
-     %0 =  %A, %B { operation = "ADD" } :
+     %0 =  %A, %B { opType = "ADD" } :
       (!gpu.mma_matrix<16x16xf16, "COp">, !gpu.mma_matrix<16x16xf16, "COp">)
       -> !gpu.mma_matrix<16x16xf16, "COp">
     ```
   }];
 
   let arguments = (ins Variadic<GPU_MMAMatrix>:$args,
-                       MMAElementWiseAttr:$operation);
+                       MMAElementWiseAttr:$opType);
 
   let results = (outs GPU_MMAMatrix:$res);
 
   let extraClassDeclaration = [{
     gpu::MMAMatrixType getType() {
       return res().getType().cast<gpu::MMAMatrixType>();
     }
   }];
 
   let assemblyFormat = [{
-    $operation $args attr-dict `:` functional-type($args, $res)
+    $opType $args attr-dict `:` functional-type($args, $res)
   }];
 }
 
 #endif // GPU_OPS

mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td

def OpenACC_Dialect : Dialect {
   let description = [{
     This dialect models the construct from the OpenACC 3.1 directive language.
   }];
 
   let useDefaultAttributePrinterParser = 1;
   let cppNamespace = "::mlir::acc";
 
   // TODO: Flip to _Prefixed.
-  let emitAccessorPrefix = kEmitAccessorPrefix_Raw;
+  let emitAccessorPrefix = kEmitAccessorPrefix_Both;
 }
 
 // AccCommon requires definition of OpenACC_Dialect.
 include "mlir/Dialect/OpenACC/AccCommon.td"
 
 // Base class for OpenACC dialect ops.
 class OpenACC_Op<string mnemonic, list<Trait> traits = []> :
   Op<OpenACC_Dialect, mnemonic, traits>;
let arguments = (ins OptionalAttr<I64Attr>:$collapse,
                        Variadic<AnyType>:$reductionOperands,
                        DefaultValuedAttr<I64Attr, "0">:$exec_mapping);
 
   let results = (outs Variadic<AnyType>:$results);
 
   let regions = (region AnyRegion:$region);
 
   let extraClassDeclaration = [{
-    static StringRef getCollapseAttrName() { return "collapse"; }
-    static StringRef getSeqAttrName() { return "seq"; }
-    static StringRef getIndependentAttrName() { return "independent"; }
-    static StringRef getAutoAttrName() { return "auto"; }
-    static StringRef getExecutionMappingAttrName() { return "exec_mapping"; }
+    static StringRef getCollapseAttrStrName() { return "collapse"; }
+    static StringRef getSeqAttrStrName() { return "seq"; }
+    static StringRef getIndependentAttrStrName() { return "independent"; }
+    static StringRef getAutoAttrStrName() { return "auto"; }
+    static StringRef getExecutionMappingAttrStrName() { return "exec_mapping"; }
     static StringRef getGangKeyword() { return "gang"; }
     static StringRef getGangNumKeyword() { return "num"; }
     static StringRef getGangStaticKeyword() { return "static"; }
     static StringRef getVectorKeyword() { return "vector"; }
     static StringRef getWorkerKeyword() { return "worker"; }
     static StringRef getTileKeyword() { return "tile"; }
     static StringRef getPrivateKeyword() { return "private"; }
     static StringRef getReductionKeyword() { return "reduction"; }

mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td

 
 def OpenMP_Dialect : Dialect {
   let name = "omp";
   let cppNamespace = "::mlir::omp";
   let dependentDialects = ["::mlir::LLVM::LLVMDialect"];
   let useDefaultAttributePrinterParser = 1;
 
   // TODO: Flip to _Prefixed.
-  let emitAccessorPrefix = kEmitAccessorPrefix_Raw;
+  let emitAccessorPrefix = kEmitAccessorPrefix_Both;
 }
 
 // OmpCommon requires definition of OpenACC_Dialect.
 include "mlir/Dialect/OpenMP/OmpCommon.td"
 
 class OpenMP_Op<string mnemonic, list<Trait> traits = []> :
       Op<OpenMP_Dialect, mnemonic, traits>;
 
oilist( `reduction` `(`
               custom<AllocateAndAllocator>(
                 $allocate_vars, type($allocate_vars),
                 $allocators_vars, type($allocators_vars)
               ) `)`
           | `proc_bind` `(` custom<ClauseAttr>($proc_bind_val) `)`
     ) $region attr-dict
   }];
   let hasVerifier = 1;
-  let extraClassDeclaration = [{
-    // TODO: remove this once emitAccessorPrefix is set to
-    // kEmitAccessorPrefix_Prefixed for the dialect.
-    /// Returns the reduction variables
-    SmallVector<Value> getReductionVars() {
-      return SmallVector<Value>(reduction_vars().begin(),
-                                reduction_vars().end());
-    }
-  }];
 }
 
 def TerminatorOp : OpenMP_Op<"terminator", [Terminator, NoSideEffect]> {
   let summary = "terminator for OpenMP regions";
   let description = [{
     A terminator operation for regions that appear in the body of OpenMP
     operation.  These regions are not expected to return any value so the
     terminator takes no operands. The terminator op returns control to the
oilist( `reduction` `(`
                 $allocators_vars, type($allocators_vars)
               ) `)`
           | `nowait` $nowait
     ) $region attr-dict
   }];
 
   let hasVerifier = 1;
   let hasRegionVerifier = 1;
-
-  let extraClassDeclaration = [{
-    // TODO: remove this once emitAccessorPrefix is set to
-    // kEmitAccessorPrefix_Prefixed for the dialect.
-    /// Returns the reduction variables
-    SmallVector<Value> getReductionVars() {
-      return SmallVector<Value>(reduction_vars().begin(),
-                                reduction_vars().end());
-    }
-  }];
 }
 
 //===----------------------------------------------------------------------===//
 // 2.8.2 Single Construct
 //===----------------------------------------------------------------------===//
 
 def SingleOp : OpenMP_Op<"single", [AttrSizedOperandSegments]> {
   let summary = "single directive";
def WsLoopOp : OpenMP_Op<"wsloop", [AttrSizedOperandSegments,
   let regions = (region AnyRegion:$region);
 
   let extraClassDeclaration = [{
     /// Returns the number of loops in the worksharing-loop nest.
     unsigned getNumLoops() { return lowerBound().size(); }
 
     /// Returns the number of reduction variables.
     unsigned getNumReductionVars() { return reduction_vars().size(); }
-
-    // TODO: remove this once emitAccessorPrefix is set to
-    // kEmitAccessorPrefix_Prefixed for the dialect.
-    /// Returns the reduction variables
-    SmallVector<Value> getReductionVars() {
-      return SmallVector<Value>(reduction_vars().begin(),
-                                reduction_vars().end());
-    }
   }];
   let hasCustomAssemblyFormat = 1;
   let assemblyFormat = [{
     oilist(`linear` `(`
               custom<LinearClause>($linear_vars, type($linear_vars),
                                    $linear_step_vars) `)`
           |`schedule` `(`
               custom<ScheduleClause>(
oilist(`if` `(` $if_expr `)`
           |`num_tasks` `(` $num_tasks `:` type($num_tasks) `)`
           |`nogroup` $nogroup
     ) `for` custom<LoopControl>($region, $lowerBound, $upperBound, $step,
                                   type($step), $inclusive) attr-dict
   }];
 
   let extraClassDeclaration = [{
     /// Returns the reduction variables
-    SmallVector<Value> getReductionVars();
+    SmallVector<Value> getAllReductionVars();
     void getEffects(SmallVectorImpl<MemoryEffects::EffectInstance> &effects);
   }];
 
   let hasVerifier = 1;
 }
 
 def TaskGroupOp : OpenMP_Op<"taskgroup", [AttrSizedOperandSegments,
                             ReductionClauseInterface,
oilist(`task_reduction` `(`
                 $allocate_vars, type($allocate_vars),
                 $allocators_vars, type($allocators_vars)
               ) `)`
     ) $region attr-dict
   }];
 
   let extraClassDeclaration = [{
     /// Returns the reduction variables
-    operand_range getReductionVars() { return task_reduction_vars(); }
+    operand_range getAllReductionVars() { return task_reduction_vars(); }
   }];
 
   let hasVerifier = 1;
 
 }
 
 //===----------------------------------------------------------------------===//
 // 2.10.4 taskyield Construct

mlir/include/mlir/Dialect/OpenMP/OpenMPOpsInterfaces.td

let description = [{
     OpenMP operations that support reduction clause have this interface.
   }];
 
   let cppNamespace = "::mlir::omp";
 
   let methods = [
     InterfaceMethod<
       "Get reduction vars", "::mlir::SmallVector<::mlir::Value>",
-      "getReductionVars">,
+      "getAllReductionVars", (ins), [{}], [{
+        return $_op.getReductionVars();
+      }]>,
   ];
 }
 
 #endif // OpenMP_OPS_INTERFACES

mlir/lib/Conversion/GPUToNVVM/WmmaOpsToNvvm.cpp

LLVM::LLVMStructType destType = convertMMAToLLVMType(
         subgroupMmaElementwiseOp.getType().cast<gpu::MMAMatrixType>());
     Value matrixStruct = rewriter.create<LLVM::UndefOp>(loc, destType);
     for (size_t i = 0, e = destType.getBody().size(); i < e; ++i) {
       SmallVector<Value> extractedOperands;
       for (size_t opIdx = 0; opIdx < numOperands; opIdx++) {
         extractedOperands.push_back(rewriter.create<LLVM::ExtractValueOp>(
             loc, adaptor.getOperands()[opIdx], i));
       }
-      Value element =
-          createScalarOp(rewriter, loc, subgroupMmaElementwiseOp.operation(),
-                         extractedOperands);
+      Value element = createScalarOp(
+          rewriter, loc, subgroupMmaElementwiseOp.opType(), extractedOperands);
       matrixStruct =
           rewriter.create<LLVM::InsertValueOp>(loc, matrixStruct, element, i);
     }
     rewriter.replaceOp(subgroupMmaElementwiseOp, matrixStruct);
     return success();
   }
 };
 

mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp

LogicalResult GPUModuleConversion::matchAndRewrite(
   // Add a keyword to the module name to avoid symbolic conflict.
   std::string spvModuleName = (kSPIRVModule + moduleOp.getName()).str();
   auto spvModule = rewriter.create<spirv::ModuleOp>(
       moduleOp.getLoc(), addressingModel, *memoryModel, llvm::None,
       StringRef(spvModuleName));
 
   // Move the region from the module op into the SPIR-V module.
   Region &spvModuleRegion = spvModule.getRegion();
-  rewriter.inlineRegionBefore(moduleOp.body(), spvModuleRegion,
+  rewriter.inlineRegionBefore(moduleOp.getBodyRegion(), spvModuleRegion,
                               spvModuleRegion.begin());
   // The spv.module build method adds a block. Remove that.
   rewriter.eraseBlock(&spvModuleRegion.back());
   rewriter.eraseOp(moduleOp);
   return success();
 }
 
 //===----------------------------------------------------------------------===//

mlir/lib/Dialect/Async/IR/Async.cpp

 //===----------------------------------------------------------------------===//
 // YieldOp
 //===----------------------------------------------------------------------===//
 
 LogicalResult YieldOp::verify() {
   // Get the underlying value types from async values returned from the
   // parent `async.execute` operation.
   auto executeOp = (*this)->getParentOfType<ExecuteOp>();
-  auto types = llvm::map_range(executeOp.results(), [](const OpResult &result) {
+  auto types =
+      llvm::map_range(executeOp.bodyResults(), [](const OpResult &result) {
         return result.getType().cast<ValueType>().getValueType();
       });
 
   if (getOperandTypes() != types)
     return emitOpError("operand types do not match the types returned from "
                        "the parent ExecuteOp");
 
   return success();
 }
 
 MutableOperandRange
 YieldOp::getMutableSuccessorOperands(Optional<unsigned> index) {
   return operandsMutable();
 }
 
 //===----------------------------------------------------------------------===//
 /// ExecuteOp
 //===----------------------------------------------------------------------===//
 
 constexpr char kOperandSegmentSizesAttr[] = "operand_segment_sizes";
 
 OperandRange ExecuteOp::getSuccessorEntryOperands(Optional<unsigned> index) {
   assert(index && *index == 0 && "invalid region index");
-  return operands();
+  return bodyOperands();
 }
 
 bool ExecuteOp::areTypesCompatible(Type lhs, Type rhs) {
   const auto getValueOrTokenType = [](Type type) {
     if (auto value = type.dyn_cast<ValueType>())
       return value.getValueType();
     return type;
   };
   return getValueOrTokenType(lhs) == getValueOrTokenType(rhs);
 }
 
 void ExecuteOp::getSuccessorRegions(Optional<unsigned> index,
                                     ArrayRef<Attribute>,
                                     SmallVectorImpl<RegionSuccessor> &regions) {
   // The `body` region branch back to the parent operation.
   if (index) {
     assert(*index == 0 && "invalid region index");
-    regions.push_back(RegionSuccessor(results()));
+    regions.push_back(RegionSuccessor(bodyResults()));
     return;
   }
 
   // Otherwise the successor is the body region.
-  regions.push_back(RegionSuccessor(&body(), body().getArguments()));
+  regions.push_back(
+      RegionSuccessor(&bodyRegion(), bodyRegion().getArguments()));
 }
 
 void ExecuteOp::build(OpBuilder &builder, OperationState &result,
                       TypeRange resultTypes, ValueRange dependencies,
                       ValueRange operands, BodyBuilderFn bodyBuilder) {
 
   result.addOperands(dependencies);
   result.addOperands(operands);

 }
 
 void ExecuteOp::print(OpAsmPrinter &p) {
   // [%tokens,...]
   if (!dependencies().empty())
     p << " [" << dependencies() << "]";
 
   // (%value as %unwrapped: !async.value<!arg.type>, ...)
-  if (!operands().empty()) {
+  if (!bodyOperands().empty()) {
     p << " (";
-    Block *entry = body().empty() ? nullptr : &body().front();
-    llvm::interleaveComma(operands(), p, [&, n = 0](Value operand) mutable {
+    Block *entry = bodyRegion().empty() ? nullptr : &bodyRegion().front();
+    llvm::interleaveComma(bodyOperands(), p, [&, n = 0](Value operand) mutable {
       Value argument = entry ? entry->getArgument(n++) : Value();
       p << operand << " as " << argument << ": " << operand.getType();
     });
     p << ")";
   }
 
   // -> (!async.value<!return.type>, ...)
   p.printOptionalArrowTypeList(llvm::drop_begin(getResultTypes()));
   p.printOptionalAttrDictWithKeyword((*this)->getAttrs(),
                                      {kOperandSegmentSizesAttr});
   p << ' ';
-  p.printRegion(body(), /*printEntryBlockArgs=*/false);
+  p.printRegion(bodyRegion(), /*printEntryBlockArgs=*/false);
 }
 
 ParseResult ExecuteOp::parse(OpAsmParser &parser, OperationState &result) {
   MLIRContext *ctx = result.getContext();
 
   // Sizes of parsed variadic operands, will be updated below after parsing.
   int32_t numDependencies = 0;
 
ParseResult ExecuteOp::parse(OpAsmParser &parser, OperationState &result) {
 
   // Parse asynchronous region.
   Region *body = result.addRegion();
   return parser.parseRegion(*body, /*arguments=*/unwrappedArgs);
 }
 
 LogicalResult ExecuteOp::verifyRegions() {
   // Unwrap async.execute value operands types.
-  auto unwrappedTypes = llvm::map_range(operands(), [](Value operand) {
+  auto unwrappedTypes = llvm::map_range(bodyOperands(), [](Value operand) {
     return operand.getType().cast<ValueType>().getValueType();
   });
 
   // Verify that unwrapped argument types matches the body region arguments.
-  if (body().getArgumentTypes() != unwrappedTypes)
+  if (bodyRegion().getArgumentTypes() != unwrappedTypes)
     return emitOpError("async body region argument types do not match the "
                        "execute operation arguments types");
 
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 /// CreateGroupOp

mlir/lib/Dialect/Async/Transforms/AsyncToAsyncRuntime.cpp

 outlineExecuteOp(SymbolTable &symbolTable, ExecuteOp execute) {
   ModuleOp module = execute->getParentOfType<ModuleOp>();
 
   MLIRContext *ctx = module.getContext();
   Location loc = execute.getLoc();
 
   // Make sure that all constants will be inside the outlined async function to
   // reduce the number of function arguments.
-  cloneConstantsIntoTheRegion(execute.body());
+  cloneConstantsIntoTheRegion(execute.bodyRegion());
 
   // Collect all outlined function inputs.
   SetVector<mlir::Value> functionInputs(execute.dependencies().begin(),
                                         execute.dependencies().end());
-  functionInputs.insert(execute.operands().begin(), execute.operands().end());
-  getUsedValuesDefinedAbove(execute.body(), functionInputs);
+  functionInputs.insert(execute.bodyOperands().begin(),
+                        execute.bodyOperands().end());
+  getUsedValuesDefinedAbove(execute.bodyRegion(), functionInputs);
 
   // Collect types for the outlined function inputs and outputs.
   auto typesRange = llvm::map_range(
       functionInputs, [](Value value) { return value.getType(); });
   SmallVector<Type, 4> inputTypes(typesRange.begin(), typesRange.end());
   auto outputTypes = execute.getResultTypes();
 
   auto funcType = FunctionType::get(ctx, inputTypes, outputTypes);
   auto funcAttrs = ArrayRef<NamedAttribute>();
 
   // TODO: Derive outlined function name from the parent FuncOp (support
   // multiple nested async.execute operations).
   func::FuncOp func =
       func::FuncOp::create(loc, kAsyncFnPrefix, funcType, funcAttrs);
   symbolTable.insert(func);
 
   SymbolTable::setSymbolVisibility(func, SymbolTable::Visibility::Private);
   auto builder = ImplicitLocOpBuilder::atBlockBegin(loc, func.addEntryBlock());
 
   // Prepare for coroutine conversion by creating the body of the function.
   {
     size_t numDependencies = execute.dependencies().size();
-    size_t numOperands = execute.operands().size();
+    size_t numOperands = execute.bodyOperands().size();
 
     // Await on all dependencies before starting to execute the body region.
     for (size_t i = 0; i < numDependencies; ++i)
       builder.create<AwaitOp>(func.getArgument(i));
 
     // Await on all async value operands and unwrap the payload.
     SmallVector<Value, 4> unwrappedOperands(numOperands);
     for (size_t i = 0; i < numOperands; ++i) {
       Value operand = func.getArgument(numDependencies + i);
       unwrappedOperands[i] = builder.create<AwaitOp>(loc, operand).result();
     }
 
     // Map from function inputs defined above the execute op to the function
     // arguments.
     BlockAndValueMapping valueMapping;
     valueMapping.map(functionInputs, func.getArguments());
-    valueMapping.map(execute.body().getArguments(), unwrappedOperands);
+    valueMapping.map(execute.bodyRegion().getArguments(), unwrappedOperands);
 
     // Clone all operations from the execute operation body into the outlined
     // function body.
-    for (Operation &op : execute.body().getOps())
+    for (Operation &op : execute.bodyRegion().getOps())
       builder.clone(op, valueMapping);
   }
 
   // Adding entry/cleanup/suspend blocks.
   CoroMachinery coro = setupCoroMachinery(func);
 
   // Suspend async function at the end of an entry block, and resume it using
   // Async resume operation (execution will be resumed in a thread managed by

mlir/lib/Dialect/GPU/IR/GPUDialect.cpp

auto walkResult = module.walk([&module](LaunchFuncOp launchOp) -> WalkResult {
     // module we are currently checking.
     if (!launchOp->getParentOp() ||
         launchOp->getParentOp()->getParentOp() != module)
       return success();
 
     // Ignore launch ops with missing attributes here. The errors will be
     // reported by the verifiers of those ops.
     if (!launchOp->getAttrOfType<SymbolRefAttr>(
-            LaunchFuncOp::getKernelAttrName()))
+            LaunchFuncOp::getKernelAttrName(launchOp->getName())))
       return success();
 
     // Check that `launch_func` refers to a well-formed GPU kernel module.
     StringAttr kernelModuleName = launchOp.getKernelModuleName();
     auto kernelModule = module.lookupSymbol<GPUModuleOp>(kernelModuleName);
     if (!kernelModule)
       return launchOp.emitOpError()
              << "kernel module '" << kernelModuleName.getValue()
result.addOperands({gridSize.x, gridSize.y, gridSize.z, blockSize.x,
                       blockSize.y, blockSize.z});
   if (dynamicSharedMemorySize)
     result.addOperands(dynamicSharedMemorySize);
   result.addOperands(kernelOperands);
   auto kernelModule = kernelFunc->getParentOfType<GPUModuleOp>();
   auto kernelSymbol =
       SymbolRefAttr::get(kernelModule.getNameAttr(),
                          {SymbolRefAttr::get(kernelFunc.getNameAttr())});
-  result.addAttribute(getKernelAttrName(), kernelSymbol);
+  result.addAttribute(getKernelAttrName(result.name), kernelSymbol);
   SmallVector<int32_t, 9> segmentSizes(9, 1);
   segmentSizes.front() = asyncDependencies.size();
   segmentSizes[segmentSizes.size() - 2] = dynamicSharedMemorySize ? 1 : 0;
   segmentSizes.back() = static_cast<int32_t>(kernelOperands.size());
   result.addAttribute(getOperandSegmentSizeAttr(),
                       builder.getDenseI32ArrayAttr(segmentSizes));
 }
 
 StringAttr LaunchFuncOp::getKernelModuleName() {
   return kernel().getRootReference();
 }
 
 StringAttr LaunchFuncOp::getKernelName() { return kernel().getLeafReference(); }
 
-unsigned LaunchFuncOp::getNumKernelOperands() { return operands().size(); }
+unsigned LaunchFuncOp::getNumKernelOperands() {
+  return kernelOperands().size();
+}
 
-Value LaunchFuncOp::getKernelOperand(unsigned i) { return operands()[i]; }
+Value LaunchFuncOp::getKernelOperand(unsigned i) { return kernelOperands()[i]; }
 
 KernelDim3 LaunchFuncOp::getGridSizeOperandValues() {
   auto operands = getOperands().drop_front(asyncDependencies().size());
   return KernelDim3{operands[0], operands[1], operands[2]};
 }
 
 KernelDim3 LaunchFuncOp::getBlockSizeOperandValues() {
   auto operands = getOperands().drop_front(asyncDependencies().size());
   return KernelDim3{operands[3], operands[4], operands[5]};
 }
 
 LogicalResult LaunchFuncOp::verify() {
   auto module = (*this)->getParentOfType<ModuleOp>();
   if (!module)
     return emitOpError("expected to belong to a module");
 
   if (!module->getAttrOfType<UnitAttr>(
           GPUDialect::getContainerModuleAttrName()))
     return emitOpError("expected the closest surrounding module to have the '" +
                        GPUDialect::getContainerModuleAttrName() +
                        "' attribute");
 
-  auto kernelAttr = (*this)->getAttrOfType<SymbolRefAttr>(getKernelAttrName());
-  if (!kernelAttr)
-    return emitOpError("symbol reference attribute '" + getKernelAttrName() +
-                       "' must be specified");
-
   return success();
 }
 
 static ParseResult parseLaunchFuncOperands(
     OpAsmParser &parser,
     SmallVectorImpl<OpAsmParser::UnresolvedOperand> &argNames,
     SmallVectorImpl<Type> &argTypes) {
   if (parser.parseOptionalKeyword("args"))

mlir/lib/Dialect/GPU/Transforms/AsyncRegionRewriter.cpp

transform(executeOp.getResultTypes(), std::back_inserter(resultTypes),
             });
   transform(results, std::back_inserter(resultTypes),
             [](Value value) { return value.getType(); });
 
   // Clone executeOp with the extra results.
   OpBuilder builder(executeOp);
   auto newOp = builder.create<async::ExecuteOp>(
       executeOp.getLoc(), TypeRange{resultTypes}.drop_front() /*drop token*/,
-      executeOp.dependencies(), executeOp.operands());
+      executeOp.dependencies(), executeOp.bodyOperands());
   BlockAndValueMapping mapper;
   newOp.getRegion().getBlocks().clear();
   executeOp.getRegion().cloneInto(&newOp.getRegion(), mapper);
 
   // Replace executeOp with cloned one.
   executeOp.getOperation()->replaceAllUsesWith(
       newOp.getResults().drop_back(results.size()));
   executeOp.erase();
TypeSwitch<Operation *>(op)
                 builder.create<async::AwaitOp>(loc, asyncToken).result());
           // Set `it` after the inserted async.await ops.
           it = builder.getInsertionPoint();
         })
         .Case<async::ExecuteOp>([&](auto executeOp) {
           // Set `it` to the beginning of the region and add asyncTokens to the
           // async.execute operands.
           it = executeOp.getBody()->begin();
-          executeOp.operandsMutable().append(asyncTokens);
+          executeOp.bodyOperandsMutable().append(asyncTokens);
           SmallVector<Type, 1> tokenTypes(
               asyncTokens.size(), builder.getType<gpu::AsyncTokenType>());
           SmallVector<Location, 1> tokenLocs(asyncTokens.size(),
                                              executeOp.getLoc());
           copy(executeOp.getBody()->addArguments(tokenTypes, tokenLocs),
                std::back_inserter(tokens));
         });
 

 };
 
 // Callback for `async.execute` ops which repeats !gpu.async.token results
 // so that each of them is only used once.
 struct GpuAsyncRegionPass::SingleTokenUseCallback {
   void operator()(async::ExecuteOp executeOp) {
     // Extract !gpu.async.token results which have multiple uses.
     auto multiUseResults =
-        llvm::make_filter_range(executeOp.results(), [](OpResult result) {
+        llvm::make_filter_range(executeOp.bodyResults(), [](OpResult result) {
           if (result.use_empty() || result.hasOneUse())
             return false;
           auto valueType = result.getType().dyn_cast<async::ValueType>();
           return valueType &&
                  valueType.getValueType().isa<gpu::AsyncTokenType>();
         });
     if (multiUseResults.empty())
       return;
 
     // Indices within !async.execute results (i.e. without the async.token).
     SmallVector<int, 4> indices;
     transform(multiUseResults, std::back_inserter(indices),
               [](OpResult result) {
                 return result.getResultNumber() - 1; // Index without token.
               });
 
     for (auto index : indices) {
-      assert(!executeOp.results()[index].getUses().empty());
+      assert(!executeOp.bodyResults()[index].getUses().empty());
       // Repeat async.yield token result, one for each use after the first one.
-      auto uses = llvm::drop_begin(executeOp.results()[index].getUses());
+      auto uses = llvm::drop_begin(executeOp.bodyResults()[index].getUses());
       auto count = std::distance(uses.begin(), uses.end());
       auto yieldOp = cast<async::YieldOp>(executeOp.getBody()->getTerminator());
       SmallVector<Value, 4> operands(count, yieldOp.getOperand(index));
       executeOp = addExecuteResults(executeOp, operands);
       // Update 'uses' to refer to the new executeOp.
-      uses = llvm::drop_begin(executeOp.results()[index].getUses());
-      auto results = executeOp.results().take_back(count);
+      uses = llvm::drop_begin(executeOp.bodyResults()[index].getUses());
+      auto results = executeOp.bodyResults().take_back(count);
       for (auto pair : llvm::zip(uses, results))
         std::get<0>(pair).set(std::get<1>(pair));
     }
   }
 };
 
 // Replaces synchronous GPU ops in the op's region with asynchronous ones and
 // inserts the necessary synchronization (as gpu.wait ops). Assumes sequential

mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp

if (failed(parseOperandList(parser, LoopOp::getPrivateKeyword(),
     return failure();
 
   // reduction()?
   if (failed(parseOperandList(parser, LoopOp::getReductionKeyword(),
                               reductionOperands, operandTypes, result)))
     return failure();
 
   if (executionMapping != acc::OpenACCExecMapping::NONE)
-    result.addAttribute(LoopOp::getExecutionMappingAttrName(),
+    result.addAttribute(LoopOp::getExecutionMappingAttrStrName(),
                         builder.getI64IntegerAttr(executionMapping));
 
   // Parse optional results in case there is a reduce.
   if (parser.parseOptionalArrowTypeList(result.types))
     return failure();
 
   if (failed(parseRegions<LoopOp>(parser, result)))
     return failure();
if (getNumResults() > 0)
     printer << " -> (" << getResultTypes() << ")";
 
   printer << ' ';
   printer.printRegion(region(),
                       /*printEntryBlockArgs=*/false,
                       /*printBlockTerminators=*/true);
 
   printer.printOptionalAttrDictWithKeyword(
-      (*this)->getAttrs(), {LoopOp::getExecutionMappingAttrName(),
+      (*this)->getAttrs(), {LoopOp::getExecutionMappingAttrStrName(),
                             LoopOp::getOperandSegmentSizeAttr()});
 }
 
 LogicalResult acc::LoopOp::verify() {
   // auto, independent and seq attribute are mutually exclusive.
   if ((auto_() && (independent() || seq())) || (independent() && seq())) {
-    return emitError("only one of " + acc::LoopOp::getAutoAttrName() + ", " +
-                     acc::LoopOp::getIndependentAttrName() + ", " +
-                     acc::LoopOp::getSeqAttrName() +
+    return emitError("only one of " + acc::LoopOp::getAutoAttrStrName() + ", " +
+                     acc::LoopOp::getIndependentAttrStrName() + ", " +
+                     acc::LoopOp::getSeqAttrStrName() +
                      " can be present at the same time");
   }
 
   // Gang, worker and vector are incompatible with seq.
   if (seq() && exec_mapping() != OpenACCExecMapping::NONE)
     return emitError("gang, worker or vector cannot appear with the seq attr");
 
   // Check non-empty body().

mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp

 
 LogicalResult ReductionOp::verify() {
   auto *op = (*this)->getParentWithTrait<ReductionClauseInterface::Trait>();
   if (!op)
     return emitOpError() << "must be used within an operation supporting "
                             "reduction clause interface";
   while (op) {
     for (const auto &var :
-         cast<ReductionClauseInterface>(op).getReductionVars())
+         cast<ReductionClauseInterface>(op).getAllReductionVars())
       if (var == accumulator())
         return success();
     op = op->getParentWithTrait<ReductionClauseInterface::Trait>();
   }
   return emitOpError() << "the accumulator is not used by the parent";
 }
 
 //===----------------------------------------------------------------------===//

 LogicalResult TaskGroupOp::verify() {
   return verifyReductionVarList(*this, task_reductions(),
                                 task_reduction_vars());
 }
 
 //===----------------------------------------------------------------------===//
 // TaskLoopOp
 //===----------------------------------------------------------------------===//
-SmallVector<Value> TaskLoopOp::getReductionVars() {
+SmallVector<Value> TaskLoopOp::getAllReductionVars() {
   SmallVector<Value> allReductionNvars(in_reduction_vars().begin(),
                                        in_reduction_vars().end());
   allReductionNvars.insert(allReductionNvars.end(), reduction_vars().begin(),
                            reduction_vars().end());
   return allReductionNvars;
 }
 
 LogicalResult TaskLoopOp::verify() {

mlir/test/Dialect/Async/ops.mlir

func.func @empty_async_execute() -> !async.token {
 
   // CHECK: return %token : !async.token
   return %token : !async.token
 }
 
 // CHECK-LABEL: @return_async_value
 func.func @return_async_value() -> !async.value<f32> {
   // CHECK: async.execute -> !async.value<f32>
-  %token, %results = async.execute -> !async.value<f32> {
+  %token, %bodyResults = async.execute -> !async.value<f32> {
     %cst = arith.constant 1.000000e+00 : f32
     async.yield %cst : f32
   }
 
-  // CHECK: return %results : !async.value<f32>
-  return %results : !async.value<f32>
+  // CHECK: return %bodyResults : !async.value<f32>
+  return %bodyResults : !async.value<f32>
 }
 
 // CHECK-LABEL: @return_captured_value
 func.func @return_captured_value() -> !async.token {
   %cst = arith.constant 1.000000e+00 : f32
   // CHECK: async.execute -> !async.value<f32>
   %token, %results = async.execute -> !async.value<f32> {
     async.yield %cst : f32
   }
 
   // CHECK: return %token : !async.token
   return %token : !async.token
 }
 
 // CHECK-LABEL: @return_async_values
 func.func @return_async_values() -> (!async.value<f32>, !async.value<f32>) {
-  %token, %results:2 = async.execute -> (!async.value<f32>, !async.value<f32>) {
+  %token, %bodyResults:2 = async.execute -> (!async.value<f32>, !async.value<f32>) {
     %cst1 = arith.constant 1.000000e+00 : f32
     %cst2 = arith.constant 2.000000e+00 : f32
     async.yield %cst1, %cst2 : f32, f32
   }
 
-  // CHECK: return %results#0, %results#1 : !async.value<f32>, !async.value<f32>
-  return %results#0, %results#1 : !async.value<f32>, !async.value<f32>
+  // CHECK: return %bodyResults#0, %bodyResults#1 : !async.value<f32>, !async.value<f32>
+  return %bodyResults#0, %bodyResults#1 : !async.value<f32>, !async.value<f32>
 }
 
 // CHECK-LABEL: @async_token_dependencies
 func.func @async_token_dependencies(%arg0: !async.token) -> !async.token {
   // CHECK: async.execute [%arg0]
   %token = async.execute [%arg0] {
     async.yield
   }

mlir/test/Dialect/GPU/invalid.mlir

^bb(%lhs : f32, %rhs : f32):
     "test.finish" () : () -> ()
   } : (f32) -> (f32)
   return
 }
 
 // -----
 
 func.func @shuffle_mismatching_type(%arg0 : f32, %arg1 : i32, %arg2 : i32) {
-  // expected-error@+1 {{op failed to verify that all of {value, result} have same type}}
+  // expected-error@+1 {{op failed to verify that all of {value, shuffleResult} have same type}}
   %shfl, %pred = "gpu.shuffle"(%arg0, %arg1, %arg2) { mode = #gpu<shuffle_mode xor> } : (f32, i32, i32) -> (i32, i1)
   return
 }
 
 // -----
 
 func.func @shuffle_unsupported_type(%arg0 : index, %arg1 : i32, %arg2 : i32) {
   // expected-error@+1 {{operand #0 must be i32 or f32}}

mlir/test/lib/Dialect/Vector/TestVectorTransforms.cpp

static Value warpReduction(Location loc, OpBuilder &builder, Value input,
                            CombiningKind kind, uint32_t size) {
   Value laneVal = input;
   // Parallel reduction using butterfly shuffles.
   for (uint64_t i = 1; i < size; i <<= 1) {
     Value shuffled = builder
                          .create<gpu::ShuffleOp>(loc, laneVal, i,
                                                  /*width=*/size,
                                                  /*mode=*/gpu::ShuffleMode::XOR)
-                         .result();
+                         .getShuffleResult();
     laneVal = makeArithReduction(builder, loc, kind, laneVal, shuffled);
   }
   return laneVal;
 }
 
 struct TestVectorDistribution
     : public PassWrapper<TestVectorDistribution, OperationPass<func::FuncOp>> {
   MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestVectorDistribution)