diff --git a/mlir/include/mlir/Dialect/GPU/IR/GPUBase.td b/mlir/include/mlir/Dialect/GPU/IR/GPUBase.td --- a/mlir/include/mlir/Dialect/GPU/IR/GPUBase.td +++ b/mlir/include/mlir/Dialect/GPU/IR/GPUBase.td @@ -96,15 +96,16 @@ }], "OperandRange", "getAsyncDependencies", (ins), [{}], [{ ConcreteOp op = cast(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), [{}], [{ - ::mlir::gpu::addAsyncDependency(this->getOperation(), token); + if (!::llvm::is_contained(this->getAsyncDependencies(), token)) + ::mlir::gpu::addAsyncDependency(this->getOperation(), token); }] >, InterfaceMethod<[{ diff --git a/mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h b/mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h --- a/mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h +++ b/mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h @@ -25,6 +25,7 @@ #include "mlir/Interfaces/InferIntRangeInterface.h" #include "mlir/Interfaces/InferTypeOpInterface.h" #include "mlir/Interfaces/SideEffectInterfaces.h" +#include "llvm/ADT/STLExtras.h" namespace mlir { namespace gpu { diff --git a/mlir/test/Dialect/GPU/async-region.mlir b/mlir/test/Dialect/GPU/async-region.mlir --- a/mlir/test/Dialect/GPU/async-region.mlir +++ b/mlir/test/Dialect/GPU/async-region.mlir @@ -175,7 +175,7 @@ // 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>