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

Update CDP support macros for if-target compatibility #1661

Merged
merged 3 commits into from
Jun 29, 2022

Conversation

alliepiper
Copy link
Collaborator

@alliepiper alliepiper commented Apr 14, 2022

Requires NVIDIA/cub#486.

This PR picks up where #1605 left off. It updates the CDP support in Thrust and CUB to eliminate the use of __CUDA_ARCH__.

The CUDART macros have been revamped (see relnotes below), and a new THRUST_CDP_DISPATCH abstraction centralizes the logic used to dispatch between CDP parallel and sequential implementations based on RDC state.

The testing/cuda/* tests have been updated to check both parallel CDP launches with RDC enabled and sequential fallbacks when RDC is disabled.

Pre-written Release Notes

Breaking Changes

  • Update CDP support macros for if-target compatibility #1661: Thrust’s CUDA Runtime support macros have been updated to support NV_IF_TARGET. They are now defined consistently across all host/device compilation passes. This should not affect most usages of these macros, but may require changes for some edge cases.
    • THRUST_RUNTIME_FUNCTION: Execution space annotations for functions that invoke CUDA Runtime APIs.
      • Old behavior:
        • RDC enabled: Defined to __host__ __device__
        • RDC not enabled:
          • NVCC host pass: Defined to __host__ __device__
          • NVCC device pass: Defined to __host__
      • New behavior:
        • RDC enabled: Defined to __host__ __device__
        • RDC not enabled: Defined to __host__
    • __THRUST_HAS_CUDART__: No change in behavior, but no longer used in Thrust. Provided for legacy support only. Legacy behavior:
      • RDC enabled: Defined to 1.
      • RDC not enabled:
        • NVCC host pass: Defined to 1.
        • NVCC device pass: Defined to 0.
    • THRUST_RDC_ENABLED: New macro, may be combined with NV_IF_TARGET to replace most usages of __THRUST_HAS_CUDART__. Behavior:
      • RDC enabled: Macro is defined.
      • RDC not enabled: Macro is not defined.

@alliepiper alliepiper marked this pull request as draft April 14, 2022 21:42
@alliepiper alliepiper self-assigned this Apr 14, 2022
@alliepiper
Copy link
Collaborator Author

run tests

@alliepiper
Copy link
Collaborator Author

run tests

1 similar comment
@alliepiper
Copy link
Collaborator Author

run tests

@alliepiper
Copy link
Collaborator Author

run tests

@alliepiper
Copy link
Collaborator Author

run tests

@alliepiper
Copy link
Collaborator Author

run tests

@alliepiper
Copy link
Collaborator Author

run tests

@alliepiper
Copy link
Collaborator Author

run tests

@alliepiper
Copy link
Collaborator Author

run tests

@alliepiper
Copy link
Collaborator Author

run tests

@alliepiper
Copy link
Collaborator Author

run tests

@alliepiper alliepiper force-pushed the if_target_cdp branch 2 times, most recently from 7485284 to c5727d9 Compare May 18, 2022 15:37
@alliepiper
Copy link
Collaborator Author

run tests

@alliepiper alliepiper added this to the 2.0.0 milestone May 18, 2022
@alliepiper alliepiper added the P0: must have Absolutely necessary. Critical issue, major blocker, etc. label May 18, 2022
@alliepiper
Copy link
Collaborator Author

run tests

@alliepiper
Copy link
Collaborator Author

run tests

@alliepiper alliepiper changed the title WIP if-target CDP refactor CI testing Update CDP support macros for if-target compatibilty May 19, 2022
@alliepiper alliepiper marked this pull request as ready for review May 19, 2022 14:57
@alliepiper alliepiper changed the title Update CDP support macros for if-target compatibilty Update CDP support macros for if-target compatibility May 19, 2022
@alliepiper alliepiper added the release: notes PR description contains pre-written release notes. label May 19, 2022
thrust/system/cuda/detail/util.h Outdated Show resolved Hide resolved
thrust/system/cuda/detail/util.h Outdated Show resolved Hide resolved
@alliepiper
Copy link
Collaborator Author

run tests

@alliepiper
Copy link
Collaborator Author

run tests

@alliepiper alliepiper merged commit 967924e into NVIDIA:main Jun 29, 2022
@alliepiper alliepiper deleted the if_target_cdp branch June 29, 2022 16:12
Sign up for free to subscribe to this conversation on GitHub. Already have an account? Sign in.
Labels
P0: must have Absolutely necessary. Critical issue, major blocker, etc. release: breaking change Include in "Breaking Changes" section of release notes. release: notes PR description contains pre-written release notes.
Projects
None yet
Development

Successfully merging this pull request may close these issues.

2 participants