Skip to content
This repository has been archived by the owner on Mar 21, 2024. It is now read-only.

Evaluate CUDA_CUB_RET_IF_FAIL macro argument only once #1264

Merged
merged 2 commits into from
Aug 28, 2020
Merged

Evaluate CUDA_CUB_RET_IF_FAIL macro argument only once #1264

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
4d567c8
Evaluate CUDA_CUB_RET_IF_FAIL macro argument only once
jlowe Aug 26, 2020
6fcf17c
Use a scope block for CUDA_CUB_RET_IF_FAIL macro
jlowe Aug 26, 2020
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 4 additions & 1 deletion thrust/system/cuda/detail/core/util.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -652,7 +652,10 @@ namespace core {
}

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

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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

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

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

// uninitialized
// -------
Expand Down