[Question]: Why does nvshmem_fence default to using nvshmemi_fence<NVSHMEMI_THREADGROUP_THREAD>()?

[kwu130]

Question

Does using NVSHMEMI_THREADGROUP_THREAD as the default scope cause excessive redundant work? Specifically, when nvshmem_fence() is called from a warp or block, all threads execute nvshmemi_ibgda_fence(), each seeing index_in_scope == 0 and scope_size == 1, and thus redundantly iterating over all DCIs and RC QPs to issue ibgda_quiet(qp) calls. Could this lead to significant performance overhead?

[akhillanger]

Yes, that is correct. It will lead to each thread iterating over all the QPs. You may want to look at the new QP APIs in NVSHMEM 3.5 that support fence API at warp/block level.

[kwu130]

Yes, that is correct. It will lead to each thread iterating over all the QPs. You may want to look at the new QP APIs in NVSHMEM 3.5 that support fence API at warp/block level.

Got it, thanks for your reply.