Evaluate CUDA_CUB_RET_IF_FAIL macro argument only once

[jlowe]

This updates the CUDA_CUB_RET_IF_FAIL macro to only evaluate its argument once. Without this, it can silently suppress CUDA errors in many places in the Thrust code that call it like this: CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());.

The macro calls cub::Debug which in some versions of cub will unconditionally clear the CUDA error. Therefore when it evaluates cudaPeekAtLastError() the second time as the return argument, it will return cudaSuccess because the last error was just cleared. Thus any pending CUDA error before this macro was called with cudaPeekAtLastErrror() as the argument will be silently dropped.

[alliepiper]

Looks good to me, aside from a minor nitpick. Can you update this patch to remove the do/while no-op? I'll start our CI afterwards.

[alliepiper]

Nitpick: The do..while isn't necessary here, this can just be

{ \
    auto const error = (e);     \
    if (cub::Debug(error, __FILE__, __LINE__)) \
      return error; \
}

[vuule]

@allisonvacanti There are cases where using a scope block causes issues, please see https://stackoverflow.com/a/1067238

[dkolsen-pgi]

For that to work, you can't have the semicolon after the while (0). With that semicolon there, it has the same problem as if it were a block scope.

[alliepiper]

To update this thread with a slack convo, we don't want that behavior here since the old version didn't require a trailing semicolon. Using the do/while(0) pattern (without the semicolon) risks breaking other code.

[alliepiper]

Started CI under DVS CL 28995185.

[alliepiper]

One DVS builder failed after running out of disk space, but the rest look good.

Just waiting for my local testing script to finish up and then this should be good to go.

[alliepiper]

All set.