[mlir][GPU] Prevent adding duplicate async tokens

[krzysz00]

Summary

If, in the GPU async transformation, the operation being given an
async dependency already depended on the token in question, we
would add duplicate tokens, creating issues in GPU to LLVM lowering.

To resolve this issue, add a check to addAsyncDependency() to ensure
that duplicate tokens are not present in the token list.

(I'm open to a different approach here, this is just what I went with
initially)

[Mogball]

mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h

29–30

[krzysz00]

mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h

29–30

I noticed a convention of separating MLIR and LLVM includes in some other files, is there not one?

[krzysz00]

Remove newline per review feedback

[krzysz00]

Somehow, getAsyncDependencies() didn't get converted to prefixed form.

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

def GPU_AsyncOpInterface : OpInterface<"AsyncOpInterface"> {
   let cppNamespace = "::mlir::gpu";
 
   let methods = [
     InterfaceMethod<[{
         Query the operands that represent async dependency tokens.
       }],
       "OperandRange", "getAsyncDependencies", (ins), [{}], [{
         ConcreteOp op = cast<ConcreteOp>(this->getOperation());
-        return op.asyncDependencies();
+        return op.getAsyncDependencies();
       }]
     >,
     InterfaceMethod<[{
-        Adds a new token to the list of async dependencies.
+        Adds a new token to the list of async dependencies if it is not already there.
       }],
       "void", "addAsyncDependency", (ins "Value":$token),
       [{}], [{
+        if (!::llvm::is_contained(this->getAsyncDependencies(), token))
           ::mlir::gpu::addAsyncDependency(this->getOperation(), token);
       }]
     >,
     InterfaceMethod<[{
         Query the result that represents the async token to depend on.
       }],
       "Value", "getAsyncToken"
     >
   ];

mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h

 #include "mlir/IR/Dialect.h"
 #include "mlir/IR/FunctionInterfaces.h"
 #include "mlir/IR/OpDefinition.h"
 #include "mlir/IR/OpImplementation.h"
 #include "mlir/IR/SymbolTable.h"
 #include "mlir/Interfaces/InferIntRangeInterface.h"
 #include "mlir/Interfaces/InferTypeOpInterface.h"
 #include "mlir/Interfaces/SideEffectInterfaces.h"
+#include "llvm/ADT/STLExtras.h"
 
 namespace mlir {
 namespace gpu {
 
 /// Utility class for the GPU dialect to represent triples of `Value`s
 /// accessible through `.x`, `.y`, and `.z` similarly to CUDA notation.
 struct KernelDim3 {
   Value x;
   Value y;
   Value z;

mlir/test/Dialect/GPU/async-region.mlir

func.func @async_execute_fork(%sz : index) {
     return
   }
 
   // CHECK-LABEL:func @existing_tokens()
   func.func @existing_tokens() {
     // CHECK: %[[t0:.*]] = gpu.wait async
     // CHECK-NOT: [{{.*}}]
     %t0 = gpu.wait async
-    // CHECK: %[[t1:.*]] = gpu.wait async [%[[t0]], %[[t0]]]
+    // CHECK: %[[t1:.*]] = gpu.wait async [%[[t0]]]
     %t1 = gpu.wait async [%t0]
     // CHECK: %[[m:.*]], %[[t2:.*]] = gpu.alloc async [%[[t1]], %[[t0]]] ()
     %0 = gpu.alloc [%t0] () : memref<7xf32>
     // CHECK: %[[t3:.*]] = gpu.dealloc async [%[[t2]]] %[[m]]
     %t2 = gpu.dealloc async %0 : memref<7xf32>
     // CHECK: gpu.wait [%[[t3]]]
     gpu.wait
     // CHECK: gpu.wait
     // CHECK-NOT: async
     // CHECK-NOT: [{{.*}}]
     gpu.wait
     return
   }
 }