Commit & Thanks

[Pranshu-Bahadur]

MatGPTQ/inference_lib/matgptq_cuda/csrc/matgptq_gemm.cu

Line 836 in 0e5bd93

cp_async_wait_all(); // TODO: Move this somewhere useful.

If you're doing pipeline commits, you need to wait_group (corresponding to a specific commit_group).
You need to cp.async.wait_group N corresponding to a specific commit group (pipeline_commit call).

As seen in in ptx isa here:

cp.async.wait_group instruction will cause executing thread to wait till only N or fewer of the most recent cp.async-groups are pending and all the prior cp.async-groups committed by the executing threads are complete. For example, when N is 0, the executing thread waits on all the prior cp.async-groups to complete. Operand N is an integer constant.

cp.async.wait_all is equivalent to :

cp.async.commit_group;
cp.async.wait_group 0;

So essentially, the wait_all calls are only synchronizing the group 0s (first pipeline_commit calls). If you did not __syncthreads(), the code would break in logic.

On that note regarding, the next statement:

MatGPTQ/inference_lib/matgptq_cuda/csrc/matgptq_gemm.cu

Line 837 in 0e5bd93

__syncthreads(); // TODO: Figure out why this is necessary?

Your comment here is a gift to me, it forced me to re-evaluate my understanding of cp.async.wait_all. I think you may be right! We dont need to syncthreads after wait_all.

An empty cp.async-group is considered to be trivially complete.

Writes performed by cp.async operations are made visible to the executing thread only after:

The completion of cp.async.wait_all or

The completion of cp.async.wait_group on the cp.async-group in which the cp.async belongs to or

mbarrier.test_wait returns True on an mbarrier object which is tracking the completion of the cp.async operation.

There is no ordering between two cp.async operations that are not synchronized with cp.async.wait_all or cp.async.wait_group or mbarrier objects.

cp.async.wait_group and cp.async.wait_all does not provide any ordering and visibility guarantees for any other memory operation apart from cp.async.

Thank you so much, I was also very unclear on this until now!