Script.clc.query_response

Script.clc.query_response

Script.clc.query_response(response)[source]

Query the response from a cluster launch control try_cancel operation.

This instruction decodes the opaque 16-byte response written by the try_cancel instruction. It extracts two pieces of information: 1. Whether the cancellation request succeeded (is_canceled) 2. If successful, the CTA ID (x, y, z coordinates) of the first CTA in the canceled cluster

The response should be loaded from shared memory after the mbarrier used in try_cancel has signaled completion of the asynchronous operation.

If the cancellation request failed, the CTA ID coordinates in the returned Dim3 are undefined and should not be used.

Parameters:

response (SharedTensor) – A 16-byte wide shared memory tensor containing the opaque response from a try_cancel operation. This should be the same tensor that was passed to try_cancel.

Returns:

  • is_canceled (Var) – A variable (predicate/boolean) that is True if the cluster was successfully canceled, False otherwise.

  • first_cta_id (Dim3) – If the cancellation succeeded, this contains the (x, y, z) coordinates of the first CTA in the canceled cluster. If the cancellation failed, these values are undefined.

Return type:

tuple[Var, Dim3]

Notes

  • Thread group: Can be executed by any sized thread group.

  • Hardware: Requires compute capability 10.0+ (sm_100).

  • PTX: clusterlaunchcontrol.query_cancel.is_canceled and clusterlaunchcontrol.query_cancel.get_first_ctaid

  • The behavior is undefined if called before the try_cancel operation has completed (i.e., before the associated mbarrier has signaled completion).

Contents