[mlir][NVGPU]: Fix op description of nvgpu.device_async_wait.

[yaoyuannnn]

Summary

According to the NVIDIA documentation on cp.async.wait_group
(https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-wait-group-cp-async-wait-all),
the numGroups attribute in nvgpu.device_async_wait should give an upper
bound of pending async group count (instead of a lower bound) when the
executing thread can be unblocked.

[yaoyuannnn]

This comment was removed by yaoyuannnn.

[christopherbate]

In case anyone needs an example to verify: if you push 16 groups, and set the wait to numGroups = 12, then that means you want to unblock when 12 groups or fewer are in flight (4 groups have completed).

It might be a bit clearer if a better example is provided.

[yaoyuannnn]

Updated with an example of using numGroups.

[yaoyuannnn]

Fixed a typo.

[yaoyuannnn]

Hi Christopher, can you help push the patch?

mlir/include/mlir/Dialect/NVGPU/IR/NVGPU.td

 }
 
 def NVGPU_DeviceAsyncWaitOp : NVGPU_Op<"device_async_wait", []> {
   let summary = "Wait for async gpu ops to complete.";
   let description = [{
     The `nvgpu.device_async_wait` op will block the execution thread until the group
     associated with the source token is fully completed.
 
-    The optional `$numGroup` attribute gives a lower bound of the number of
-    groups uncompleted when the wait can unblock the thread.
+    The optional `$numGroups` attribute gives an upper bound of the number of
+    groups uncompleted when the wait can unblock the thread. For example,  if
+    16 async groups are pushe and `$numGroups` is set to 12, then the thread
+    will unblock when 12 groups or fewer are in flight (4 groups have
+    completed).
 
     Example:
 
     ```mlir
     nvgpu.device_async_wait %0
     ```
   }];
   let arguments = (ins NVGPU_DeviceAsyncToken:$asyncDependencies,
                        OptionalAttr<I32Attr>:$numGroups);
   let assemblyFormat = [{
     $asyncDependencies attr-dict
   }];
 }
 
 #endif // NVGPU