Fix a race condition bug caused by missing nullptr check

Summary:
This is a very tricky race condition bug that was root caused to the nullptr dereference of kernelFlag.

We are missing nullptr check in the KernelTestHostAbort() kernel, which caused illegal memory access during the distributed test of allgatherP on AMD:
`[1,8]<stdout>:Failed: Cuda error fbcode/comms/ctran/tests/CtranDistAllgatherPTests.cc:59 'an illegal memory access was encountered'
[1,8]<stdout>:W1119 21:46:34.875350 952795 CtranMapperRegMem.cc:197] rtptest042:952795:952795 [0][main] CTRAN WARN Total 2/2 remaining segments are still in RegCache at destroy time. `
full results see P2047923648

For Nvidia only tests, this issue didn't happen because the `abort` fileld in CtranComm [is initialized as `false`](https://fburl.com/code/0fybwtgz)
shmDevState.enableCTancellableWaits inherits this value, so ctran::utils::loadInt(flag) is short-circuited.

For NV/AMD compatible UT of allgatherP, we use the McclComm obj to initialize CtranComm. It by default enables cancellable waits, and thus exposed the dereference of flag.

The failure was very undeterministic because it relys on the race of "step" and  "val" on this line https://fburl.com/code/l24vhay5.

The fix is simple, if flag is empty, we assume the user do not need cancellable waits, and  we can simply return false in KernelTestHostAbort.

Reviewed By: cenzhaometa

Differential Revision: D87946459

fbshipit-source-id: ce4cc441768d5fbbdcff053fa9790fef0ace73b0

Author: dmwu

comms/ctran/algos/common/GpeKernelDev.cuh

@@ -45,7 +45,8 @@ static inline __device__ bool KernelTestHostAbort(volatile int* flag) {
   // local), or user active aborts/timeouts will put CtranComm in such state, if
   // abort is enabled for CtranComm.
   return shmDevState.enableCancellableWaits &&
-      (kernelDoAbort || ctran::utils::loadInt(flag) == KERNEL_HOST_ABORT);
+      (kernelDoAbort ||
+       (flag && ctran::utils::loadInt(flag) == KERNEL_HOST_ABORT));
 }
 
 static inline __device__ bool KernelTestHostAbortBlock(volatile int* flag) {