From 64e7a0689ae0e3677f5926ef6d4836c509443e63 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Thu, 18 Jul 2024 14:13:31 +0200 Subject: [PATCH 01/44] Clarify PtxVersion and SmVersion (#2004) --- cub/cub/util_device.cuh | 11 ++++------- 1 file changed, 4 insertions(+), 7 deletions(-) diff --git a/cub/cub/util_device.cuh b/cub/cub/util_device.cuh index 2998608d567..3fc4ad5bb6a 100644 --- a/cub/cub/util_device.cuh +++ b/cub/cub/util_device.cuh @@ -383,10 +383,9 @@ struct SmVersionCacheTag {}; /** - * \brief Retrieves the PTX version that will be used on \p device (major * 100 + minor * 10). + * \brief Retrieves the PTX virtual architecture that will be used on \p device (major * 100 + minor * 10). * * \note This function may cache the result internally. - * * \note This function is thread safe. */ _CCCL_HOST inline cudaError_t PtxVersion(int& ptx_version, int device) @@ -408,10 +407,9 @@ _CCCL_HOST inline cudaError_t PtxVersion(int& ptx_version, int device) } /** - * \brief Retrieves the PTX version that will be used on the current device (major * 100 + minor * 10). + * \brief Retrieves the PTX virtual architecture that will be used on the current device (major * 100 + minor * 10). * * \note This function may cache the result internally. - * * \note This function is thread safe. */ CUB_RUNTIME_FUNCTION inline cudaError_t PtxVersion(int& ptx_version) @@ -438,7 +436,7 @@ CUB_RUNTIME_FUNCTION inline cudaError_t PtxVersion(int& ptx_version) } /** - * \brief Retrieves the SM version of \p device (major * 100 + minor * 10) + * \brief Retrieves the SM version (i.e. compute capability) of \p device (major * 100 + minor * 10) */ CUB_RUNTIME_FUNCTION inline cudaError_t SmVersionUncached(int& sm_version, int device = CurrentDevice()) { @@ -464,10 +462,9 @@ CUB_RUNTIME_FUNCTION inline cudaError_t SmVersionUncached(int& sm_version, int d } /** - * \brief Retrieves the SM version of \p device (major * 100 + minor * 10) + * \brief Retrieves the SM version (i.e. compute capability) of \p device (major * 100 + minor * 10). * * \note This function may cache the result internally. - * * \note This function is thread safe. */ CUB_RUNTIME_FUNCTION inline cudaError_t SmVersion(int& sm_version, int device = CurrentDevice()) From 87d08495e2be44e718b1cc51785c434f6b6836d6 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Thu, 18 Jul 2024 18:13:30 +0200 Subject: [PATCH 02/44] Refactor CUB util_device (#1948) --- cub/cub/util_device.cuh | 44 ++++++++++++----------------------------- 1 file changed, 13 insertions(+), 31 deletions(-) diff --git a/cub/cub/util_device.cuh b/cub/cub/util_device.cuh index 3fc4ad5bb6a..b9be58d1bfc 100644 --- a/cub/cub/util_device.cuh +++ b/cub/cub/util_device.cuh @@ -155,9 +155,10 @@ CUB_RUNTIME_FUNCTION inline int DeviceCountUncached() /** * \brief Cache for an arbitrary value produced by a nullary function. + * deprecated [Since 2.6.0] */ template -struct ValueCache +struct CUB_DEPRECATED ValueCache { T const value; @@ -170,13 +171,11 @@ struct ValueCache {} }; -// Host code, only safely usable in C++11 or newer, where thread-safe -// initialization of static locals is guaranteed. This is a separate function -// to avoid defining a local static in a host/device function. +// Host code. This is a separate function to avoid defining a local static in a host/device function. _CCCL_HOST inline int DeviceCountCachedValue() { - static ValueCache cache; - return cache.value; + static int count = DeviceCountUncached(); + return count; } /** @@ -211,7 +210,7 @@ struct PerDeviceAttributeCache // Each entry starts in the `DeviceEntryEmpty` state, then proceeds to the // `DeviceEntryInitializing` state, and then proceeds to the // `DeviceEntryReady` state. These are the only state transitions allowed; - // e.g. a linear sequence of transitions. + // i.e. a linear sequence of transitions. enum DeviceEntryStatus { DeviceEntryEmpty = 0, @@ -372,7 +371,6 @@ _CCCL_HOST inline cudaError_t PtxVersionUncached(int& ptx_version, int device) template _CCCL_HOST inline PerDeviceAttributeCache& GetPerDeviceAttributeCache() { - // C++11 guarantees that initialization of static locals is thread safe. static PerDeviceAttributeCache cache; return cache; } @@ -391,8 +389,7 @@ struct SmVersionCacheTag _CCCL_HOST inline cudaError_t PtxVersion(int& ptx_version, int device) { auto const payload = GetPerDeviceAttributeCache()( - // If this call fails, then we get the error code back in the payload, - // which we check with `CubDebug` below. + // If this call fails, then we get the error code back in the payload, which we check with `CubDebug` below. [=](int& pv) { return PtxVersionUncached(pv, device); }, @@ -415,23 +412,10 @@ _CCCL_HOST inline cudaError_t PtxVersion(int& ptx_version, int device) CUB_RUNTIME_FUNCTION inline cudaError_t PtxVersion(int& ptx_version) { cudaError_t result = cudaErrorUnknown; - NV_IF_TARGET( - NV_IS_HOST, - (auto const device = CurrentDevice(); - auto const payload = GetPerDeviceAttributeCache()( - // If this call fails, then we get the error code back in the payload, - // which we check with `CubDebug` below. - [=](int& pv) { - return PtxVersionUncached(pv, device); - }, - device); - - if (!CubDebug(payload.error)) { ptx_version = payload.attribute; } - - result = payload.error;), - ( // NV_IS_DEVICE: - result = PtxVersionUncached(ptx_version);)); - + NV_IF_TARGET(NV_IS_HOST, + (result = PtxVersion(ptx_version, CurrentDevice());), + ( // NV_IS_DEVICE: + result = PtxVersionUncached(ptx_version);)); return result; } @@ -474,8 +458,7 @@ CUB_RUNTIME_FUNCTION inline cudaError_t SmVersion(int& sm_version, int device = NV_IF_TARGET( NV_IS_HOST, (auto const payload = GetPerDeviceAttributeCache()( - // If this call fails, then we get the error code back in - // the payload, which we check with `CubDebug` below. + // If this call fails, then we get the error code back in the payload, which we check with `CubDebug` below. [=](int& pv) { return SmVersionUncached(pv, device); }, @@ -562,9 +545,8 @@ CUB_RUNTIME_FUNCTION inline cudaError_t DebugSyncStream(cudaStream_t stream) CUB_RUNTIME_FUNCTION inline cudaError_t HasUVA(bool& has_uva) { has_uva = false; - cudaError_t error = cudaSuccess; int device = -1; - error = CubDebug(cudaGetDevice(&device)); + cudaError_t error = CubDebug(cudaGetDevice(&device)); if (cudaSuccess != error) { return error; From 92b4b0b607a0c2825a21705c8dae4cff4159ad96 Mon Sep 17 00:00:00 2001 From: Eric Niebler Date: Thu, 18 Jul 2024 10:37:12 -0700 Subject: [PATCH 03/44] fix some typos in `` (#2003) --- libcudacxx/include/cuda/stream_ref | 24 ++++++++++++------------ 1 file changed, 12 insertions(+), 12 deletions(-) diff --git a/libcudacxx/include/cuda/stream_ref b/libcudacxx/include/cuda/stream_ref index 9bb23d3e2ef..4cb53bbbc08 100644 --- a/libcudacxx/include/cuda/stream_ref +++ b/libcudacxx/include/cuda/stream_ref @@ -69,7 +69,7 @@ public: using value_type = ::cudaStream_t; /** - * \brief Constructs a `stream_view` of the "default" CUDA stream. + * \brief Constructs a `stream_ref` of the "default" CUDA stream. * * For behavior of the default stream, * \see @@ -79,11 +79,11 @@ public: stream_ref() = default; /** - * \brief Constructs a `stream_view` from a `cudaStream_t` handle. + * \brief Constructs a `stream_ref` from a `cudaStream_t` handle. * * This constructor provides implicit conversion from `cudaStream_t`. * - * \note: It is the callers responsibilty to ensure the `stream_view` does not + * \note: It is the callers responsibilty to ensure the `stream_ref` does not * outlive the stream identified by the `cudaStream_t` handle. * */ @@ -98,13 +98,13 @@ public: stream_ref(_CUDA_VSTD::nullptr_t) = delete; /** - * \brief Compares two `stream_view`s for equality + * \brief Compares two `stream_ref`s for equality * * \note Allows comparison with `cudaStream_t` due to implicit conversion to - * `stream_view`. + * `stream_ref`. * - * \param lhs The first `stream_view` to compare - * \param rhs The second `stream_view` to compare + * \param lhs The first `stream_ref` to compare + * \param rhs The second `stream_ref` to compare * \return true if equal, false if unequal */ _CCCL_NODISCARD_FRIEND constexpr bool operator==(const stream_ref& __lhs, const stream_ref& __rhs) noexcept @@ -113,13 +113,13 @@ public: } /** - * \brief Compares two `stream_view`s for inequality + * \brief Compares two `stream_ref`s for inequality * * \note Allows comparison with `cudaStream_t` due to implicit conversion to - * `stream_view`. + * `stream_ref`. * - * \param lhs The first `stream_view` to compare - * \param rhs The second `stream_view` to compare + * \param lhs The first `stream_ref` to compare + * \param rhs The second `stream_ref` to compare * \return true if unequal, false if equal */ _CCCL_NODISCARD_FRIEND constexpr bool operator!=(const stream_ref& __lhs, const stream_ref& __rhs) noexcept @@ -164,7 +164,7 @@ public: break; default: ::cudaGetLastError(); // Clear CUDA error state - ::cuda::__throw_cuda_error(__result, "Failed to querry stream."); + ::cuda::__throw_cuda_error(__result, "Failed to query stream."); } return true; } From 56d99db3c959bffac2cc4f0302424a3fa06c7e33 Mon Sep 17 00:00:00 2001 From: Allison Piper Date: Thu, 18 Jul 2024 15:39:19 -0400 Subject: [PATCH 04/44] Add CI slack notifications. (#1961) * Fix divide by zero in `parse-job-times.py` * Add CI slack notifications. --- .github/actions/workflow-build/action.yml | 38 ++++++++++++++ .github/actions/workflow-results/action.yml | 52 ++++++++++++++++++- .../workflow-results/parse-job-times.py | 2 +- .github/workflows/build-rapids.yml | 28 ++++++++++ .github/workflows/ci-workflow-nightly.yml | 46 ++++------------ 5 files changed, 128 insertions(+), 38 deletions(-) diff --git a/.github/actions/workflow-build/action.yml b/.github/actions/workflow-build/action.yml index 57a94989523..3842886589a 100644 --- a/.github/actions/workflow-build/action.yml +++ b/.github/actions/workflow-build/action.yml @@ -25,6 +25,15 @@ inputs: description: "Path to the matrix parser script (default if blank: build-workflow.py from action dir)" default: "" required: false + slack_token: + description: "The Slack token to use for notifications. No notifications will be sent if not provided." + required: false + slack_log: + description: "Slack channel ID for verbose notifications." + required: false + slack_alert: + description: "Slack channel ID for alert notifications." + required: false outputs: workflow: @@ -35,6 +44,20 @@ runs: using: "composite" steps: + - name: Send Slack log notification + if: ${{inputs.slack_token != '' && inputs.slack_log != '' }} + uses: slackapi/slack-github-action@v1.26.0 + env: + SLACK_BOT_TOKEN: ${{ inputs.slack_token }} + WORKFLOW_TYPE: ${{ github.workflow }} # nightly, weekly, pr, etc. + SUMMARY_URL: https://github.com/${{github.repository}}/actions/runs/${{github.run_id}} + with: + channel-id: ${{ inputs.slack_log }} + slack-message: | + Workflow '${{ env.WORKFLOW_TYPE }}' starting... + + Details: ${{ env.SUMMARY_URL }} + - name: Inspect changes if: ${{ inputs.inspect_changes_script != '' && inputs.inspect_changes_base_sha != '' }} id: inspect-changes @@ -99,3 +122,18 @@ runs: name: workflow path: workflow/ compression-level: 0 + + - name: Send Slack error notification + if: ${{ failure() && inputs.slack_token != '' && (inputs.slack_alert != '' || inputs.slack_log != '') }} + uses: slackapi/slack-github-action@v1.26.0 + env: + SLACK_BOT_TOKEN: ${{ inputs.slack_token }} + WORKFLOW_TYPE: ${{ github.workflow }} # nightly, weekly, pr, etc. + SUMMARY_URL: https://github.com/${{github.repository}}/actions/runs/${{github.run_id}} + CHANNEL_SEP: ${{ (inputs.slack_log != '' && inputs.slack_alert != '') && ',' || ''}} + with: + channel-id: '${{ inputs.slack_log }}${{env.CHANNEL_SEP}}${{ inputs.slack_alert }}' + slack-message: | + Workflow '${{ env.WORKFLOW_TYPE }}' encountered an error while preparing to run. + + Details: ${{ env.SUMMARY_URL }} diff --git a/.github/actions/workflow-results/action.yml b/.github/actions/workflow-results/action.yml index 3f5bc92afad..f14d6d496e8 100644 --- a/.github/actions/workflow-results/action.yml +++ b/.github/actions/workflow-results/action.yml @@ -8,11 +8,20 @@ inputs: pr_number: description: "The PR number to comment on, if applicable. No comment will be made if not provided." required: false + slack_token: + description: "The Slack token to use for notifications. No notifications will be sent if not provided." + required: false + slack_log: + description: "Slack channel ID for verbose notifications." + required: false + slack_alert: + description: "Slack channel ID for alert notifications." + required: false outputs: success: description: "Whether any jobs failed." - value: ${{ steps.check-dispatch.outputs.success }} + value: ${{ steps.check-success.outputs.success }} runs: using: "composite" @@ -112,6 +121,10 @@ runs: printf "SUMMARY=%s\n" "$(cat final_summary.md | url_encode_string)" | tee -a "${GITHUB_OUTPUT}" echo "::endgroup::" + echo "::group::GHA Output: EXEC_SUMMARY" + printf "EXEC_SUMMARY=%s\n" "$(cat execution/heading.txt)" | tee -a "${GITHUB_OUTPUT}" + echo "::endgroup::" + cp final_summary.md ${GITHUB_STEP_SUMMARY} - name: Comment on PR @@ -140,7 +153,7 @@ runs: }); - name: Check for job success - id: check-dispatch + id: check-success shell: bash --noprofile --norc -euo pipefail {0} run: | echo "::group::Checking for success artifacts" @@ -162,3 +175,38 @@ runs: fi echo "success=true" >> "${GITHUB_OUTPUT}" + + - name: Send Slack log notification + if: ${{ always() && inputs.slack_token != '' && inputs.slack_log != '' }} + uses: slackapi/slack-github-action@v1.26.0 + env: + SLACK_BOT_TOKEN: ${{ inputs.slack_token }} + WORKFLOW_TYPE: ${{ github.workflow }} # nightly, weekly, pr, etc. + STATUS: ${{ steps.check-success.outcome }} + EXEC_SUMMARY: ${{ steps.final-summary.outputs.EXEC_SUMMARY }} + SUMMARY_URL: https://github.com/${{github.repository}}/actions/runs/${{github.run_id}} + with: + channel-id: ${{ inputs.slack_log }} + slack-message: | + Workflow '${{ env.WORKFLOW_TYPE }}' has finished with status `${{ env.STATUS }}`: + + ${{ env.EXEC_SUMMARY }} + + Details: ${{ env.SUMMARY_URL }} + + - name: Send Slack alert notification + if: ${{ failure() && inputs.slack_token != '' && inputs.slack_alert != '' }} + uses: slackapi/slack-github-action@v1.26.0 + env: + SLACK_BOT_TOKEN: ${{ inputs.slack_token }} + WORKFLOW_TYPE: ${{ github.workflow }} # nightly, weekly, pr, etc. + EXEC_SUMMARY: ${{ steps.final-summary.outputs.EXEC_SUMMARY }} + SUMMARY_URL: https://github.com/${{github.repository}}/actions/runs/${{github.run_id}} + with: + channel-id: ${{ inputs.slack_alert }} + slack-message: | + Workflow '${{ env.WORKFLOW_TYPE }}' has failed: + + ${{ env.EXEC_SUMMARY }} + + Details: ${{ env.SUMMARY_URL }} diff --git a/.github/actions/workflow-results/parse-job-times.py b/.github/actions/workflow-results/parse-job-times.py index b90bd227a30..b30d585a0a6 100755 --- a/.github/actions/workflow-results/parse-job-times.py +++ b/.github/actions/workflow-results/parse-job-times.py @@ -120,7 +120,7 @@ def main(): for id, stats in result.items(): job_seconds = stats['job_seconds'] command_seconds = stats['command_seconds'] - overhead = (job_seconds - command_seconds) * 100 / command_seconds + overhead = (job_seconds - command_seconds) * 100 / command_seconds if command_seconds > 0 else 100 print(f"{stats['job_duration']:10} {stats['command_duration']:10} {overhead:10.0f} {stats['name']}") print("::endgroup::") diff --git a/.github/workflows/build-rapids.yml b/.github/workflows/build-rapids.yml index fe4a9697be8..9796cc81780 100644 --- a/.github/workflows/build-rapids.yml +++ b/.github/workflows/build-rapids.yml @@ -2,6 +2,15 @@ name: Build all RAPIDS repositories on: workflow_call: + inputs: + slack_token: + description: "The Slack token to use for notifications. No notifications will be sent if not provided." + required: false + type: string + slack_alert: + description: "Slack channel ID for alert notifications." + required: false + type: string jobs: check-event: @@ -156,3 +165,22 @@ jobs: --volume "$RUNNER_TEMP/ci.sh:/ci.sh" \ --volume "$RUNNER_TEMP/ci-entrypoint.sh:/ci-entrypoint.sh" \ -- /ci-entrypoint.sh ./ci/rapids/rapids-entrypoint.sh /ci.sh + + notify-failure: + name: Notify Slack of RAPIDS failure + if: ${{ failure() && inputs.slack_token != '' && inputs.slack_alert != '' }} + needs: build-rapids + runs-on: ubuntu-latest + steps: + - name: Notify + uses: slackapi/slack-github-action@v1.26.0 + env: + SLACK_BOT_TOKEN: ${{ inputs.slack_token }} + WORKFLOW_TYPE: ${{ github.workflow }} + SUMMARY_URL: https://github.com/${{github.repository}}/actions/runs/${{github.run_id}} + with: + channel-id: ${{ inputs.slack_alert }} + slack-message: | + RAPIDS build in workflow '${{ env.WORKFLOW_TYPE }}' failed. + + Details: ${{ env.SUMMARY_URL }} diff --git a/.github/workflows/ci-workflow-nightly.yml b/.github/workflows/ci-workflow-nightly.yml index c7406ef817e..11323dd2174 100644 --- a/.github/workflows/ci-workflow-nightly.yml +++ b/.github/workflows/ci-workflow-nightly.yml @@ -46,6 +46,9 @@ jobs: uses: ./.github/actions/workflow-build with: workflows: nightly + slack_token: ${{ secrets.SLACK_NOTIFIER_BOT_TOKEN }} + slack_log: ${{ secrets.SLACK_CHANNEL_CI_LOG }} + slack_alert: ${{ secrets.SLACK_CHANNEL_CI_ALERT }} dispatch-groups-linux-two-stage: name: ${{ matrix.name }} @@ -128,6 +131,11 @@ jobs: - name: Check workflow success id: check-workflow uses: ./.github/actions/workflow-results + with: + slack_token: ${{ secrets.SLACK_NOTIFIER_BOT_TOKEN }} + slack_log: ${{ secrets.SLACK_CHANNEL_CI_LOG }} + slack_alert: ${{ secrets.SLACK_CHANNEL_CI_ALERT }} + build-rapids: name: Build RAPIDS @@ -139,38 +147,6 @@ jobs: contents: read pull-requests: read uses: ./.github/workflows/build-rapids.yml - - # Check all other job statuses. This job gates branch protection checks. - ci: - name: CI - # !! Important: This job is used for branch protection checks. - # !! Need to use always() instead of !cancelled() because skipped jobs count as success - # !! for Github branch protection checks. Yes, really: by default, branch protections - # !! can be bypassed by cancelling CI. See NVIDIA/cccl#605. - if: ${{ always() }} - needs: - - verify-workflow - runs-on: ubuntu-latest - steps: - - name: Check results - run: | - status="success" - - check_result() { - name=$1 - expected=$2 - result=$3 - - echo "Checking if $name job result ('$result') is '$expected'..." - if [[ "$result" != "$expected" ]]; then - echo "$name job failed" - - status="failed" - fi - } - - check_result "verify-workflow" "success" "${{needs.verify-workflow.result}}" - - if [[ "$status" != "success" ]]; then - exit 1 - fi + with: + slack_token: ${{ secrets.SLACK_NOTIFIER_BOT_TOKEN }} + slack_alert: ${{ secrets.SLACK_CHANNEL_CI_ALERT }} From fc457b471cad86015e4788479ff28e3392234779 Mon Sep 17 00:00:00 2001 From: Allison Piper Date: Thu, 18 Jul 2024 17:01:09 -0400 Subject: [PATCH 05/44] Allow nightly workflow to be manually invoked. (#2007) [skip-rapids][skip-vdc][skip-matrix][skip-docs] --- .github/workflows/ci-workflow-nightly.yml | 1 + 1 file changed, 1 insertion(+) diff --git a/.github/workflows/ci-workflow-nightly.yml b/.github/workflows/ci-workflow-nightly.yml index 11323dd2174..5e68915e1ce 100644 --- a/.github/workflows/ci-workflow-nightly.yml +++ b/.github/workflows/ci-workflow-nightly.yml @@ -21,6 +21,7 @@ defaults: shell: bash --noprofile --norc -euo pipefail {0} on: + workflow_dispatch: schedule: - cron: '0 7 * * *' # 7AM UTC, 12AM PST, 3AM EST From eb62dc61a21a980b70b82e572365e4397f223f36 Mon Sep 17 00:00:00 2001 From: Allison Piper Date: Thu, 18 Jul 2024 18:33:06 -0400 Subject: [PATCH 06/44] Need to use a different approach to reuse secrets in reusable workflows vs. actions. (#2008) [skip-matrix][skip-vdc][skip-docs][skip-rapids] --- .github/workflows/build-rapids.yml | 17 +++++++---------- .github/workflows/ci-workflow-nightly.yml | 3 +-- 2 files changed, 8 insertions(+), 12 deletions(-) diff --git a/.github/workflows/build-rapids.yml b/.github/workflows/build-rapids.yml index 9796cc81780..4d78cc777fb 100644 --- a/.github/workflows/build-rapids.yml +++ b/.github/workflows/build-rapids.yml @@ -3,14 +3,11 @@ name: Build all RAPIDS repositories on: workflow_call: inputs: - slack_token: - description: "The Slack token to use for notifications. No notifications will be sent if not provided." + enable_slack_alerts: + description: "If true, a message will be posted to the CCCL GHA CI Alert channel if the workflow fails." required: false - type: string - slack_alert: - description: "Slack channel ID for alert notifications." - required: false - type: string + default: false + type: boolean jobs: check-event: @@ -168,18 +165,18 @@ jobs: notify-failure: name: Notify Slack of RAPIDS failure - if: ${{ failure() && inputs.slack_token != '' && inputs.slack_alert != '' }} + if: ${{ failure() && inputs.enable_slack_alerts }} needs: build-rapids runs-on: ubuntu-latest steps: - name: Notify uses: slackapi/slack-github-action@v1.26.0 env: - SLACK_BOT_TOKEN: ${{ inputs.slack_token }} + SLACK_BOT_TOKEN: ${{ secrets.SLACK_NOTIFIER_BOT_TOKEN }} WORKFLOW_TYPE: ${{ github.workflow }} SUMMARY_URL: https://github.com/${{github.repository}}/actions/runs/${{github.run_id}} with: - channel-id: ${{ inputs.slack_alert }} + channel-id: ${{ secrets.SLACK_CHANNEL_CI_ALERT }} slack-message: | RAPIDS build in workflow '${{ env.WORKFLOW_TYPE }}' failed. diff --git a/.github/workflows/ci-workflow-nightly.yml b/.github/workflows/ci-workflow-nightly.yml index 5e68915e1ce..619d2203d4f 100644 --- a/.github/workflows/ci-workflow-nightly.yml +++ b/.github/workflows/ci-workflow-nightly.yml @@ -149,5 +149,4 @@ jobs: pull-requests: read uses: ./.github/workflows/build-rapids.yml with: - slack_token: ${{ secrets.SLACK_NOTIFIER_BOT_TOKEN }} - slack_alert: ${{ secrets.SLACK_CHANNEL_CI_ALERT }} + enable_slack_alerts: true From 97e699fc3b804f7c03947fb1fd752f05c7bb0a6e Mon Sep 17 00:00:00 2001 From: Allison Piper Date: Thu, 18 Jul 2024 20:59:15 -0400 Subject: [PATCH 07/44] Enable RAPIDS builds for manually dispatched workflows. (#2009) [skip-rapids][skip-vdc][skip-matrix][skip-docs] --- .github/workflows/build-rapids.yml | 1 + 1 file changed, 1 insertion(+) diff --git a/.github/workflows/build-rapids.yml b/.github/workflows/build-rapids.yml index 4d78cc777fb..83dc9bf4f24 100644 --- a/.github/workflows/build-rapids.yml +++ b/.github/workflows/build-rapids.yml @@ -22,6 +22,7 @@ jobs: run: | [[ '${{ github.event_name }}' == 'push' && '${{ github.repository }}' == 'NVIDIA/cccl' ]] || \ [[ '${{ github.event_name }}' == 'schedule' && '${{ github.repository }}' == 'NVIDIA/cccl' ]] || \ + [[ '${{ github.event_name }}' == 'workflow_dispatch' && '${{ github.repository }}' == 'NVIDIA/cccl' ]] || \ [[ '${{ github.event_name }}' == 'pull_request' && '${{ github.repository }}' != 'NVIDIA/cccl' ]] \ && echo "ok=true" | tee -a $GITHUB_OUTPUT \ || echo "ok=false" | tee -a $GITHUB_OUTPUT; From 2ff83a2cb3daa2b2992f37345606ca5aec8c0695 Mon Sep 17 00:00:00 2001 From: Danial Javady <122740063+ZelboK@users.noreply.github.com> Date: Thu, 18 Jul 2024 21:59:31 -0400 Subject: [PATCH 08/44] Clean up complex.inl (#1655) --- thrust/thrust/detail/complex/complex.inl | 65 +----------------------- 1 file changed, 1 insertion(+), 64 deletions(-) diff --git a/thrust/thrust/detail/complex/complex.inl b/thrust/thrust/detail/complex/complex.inl index 6532f32fc67..3f4c1269312 100644 --- a/thrust/thrust/detail/complex/complex.inl +++ b/thrust/thrust/detail/complex/complex.inl @@ -1,5 +1,5 @@ /* - * Copyright 2008-2021 NVIDIA Corporation + * Copyright 2008-2024 NVIDIA Corporation * Copyright 2013 Filipe RNC Maia * * Licensed under the Apache License, Version 2.0 (the "License"); @@ -26,97 +26,44 @@ THRUST_NAMESPACE_BEGIN /* --- Constructors --- */ -#if THRUST_CPP_DIALECT < 2011 -template -_CCCL_HOST_DEVICE complex::complex() -{ - real(T()); - imag(T()); -} -#endif - template _CCCL_HOST_DEVICE complex::complex(const T& re) -#if THRUST_CPP_DIALECT >= 2011 // Initialize the storage in the member initializer list using C++ unicorn // initialization. This allows `complex` to work. : data{re, T()} {} -#else -{ - real(re); - imag(T()); -} -#endif template _CCCL_HOST_DEVICE complex::complex(const T& re, const T& im) -#if THRUST_CPP_DIALECT >= 2011 // Initialize the storage in the member initializer list using C++ unicorn // initialization. This allows `complex` to work. : data{re, im} {} -#else -{ - real(re); - imag(im); -} -#endif - -#if THRUST_CPP_DIALECT < 2011 -template -_CCCL_HOST_DEVICE complex::complex(const complex& z) -{ - real(z.real()); - imag(z.imag()); -} -#endif template template _CCCL_HOST_DEVICE complex::complex(const complex& z) -#if THRUST_CPP_DIALECT >= 2011 // Initialize the storage in the member initializer list using C++ unicorn // initialization. This allows `complex` to work. // We do a functional-style cast here to suppress conversion warnings. : data{T(z.real()), T(z.imag())} {} -#else -{ - real(T(z.real())); - imag(T(z.imag())); -} -#endif template _CCCL_HOST THRUST_STD_COMPLEX_DEVICE complex::complex(const std::complex& z) -#if THRUST_CPP_DIALECT >= 2011 // Initialize the storage in the member initializer list using C++ unicorn // initialization. This allows `complex` to work. : data{THRUST_STD_COMPLEX_REAL(z), THRUST_STD_COMPLEX_IMAG(z)} {} -#else -{ - real(THRUST_STD_COMPLEX_REAL(z)); - imag(THRUST_STD_COMPLEX_IMAG(z)); -} -#endif template template _CCCL_HOST THRUST_STD_COMPLEX_DEVICE complex::complex(const std::complex& z) -#if THRUST_CPP_DIALECT >= 2011 // Initialize the storage in the member initializer list using C++ unicorn // initialization. This allows `complex` to work. // We do a functional-style cast here to suppress conversion warnings. : data{T(THRUST_STD_COMPLEX_REAL(z)), T(THRUST_STD_COMPLEX_IMAG(z))} {} -#else -{ - real(T(THRUST_STD_COMPLEX_REAL(z))); - imag(T(THRUST_STD_COMPLEX_IMAG(z))); -} -#endif /* --- Assignment Operators --- */ @@ -128,16 +75,6 @@ _CCCL_HOST_DEVICE complex& complex::operator=(const T& re) return *this; } -#if THRUST_CPP_DIALECT < 2011 -template -_CCCL_HOST_DEVICE complex& complex::operator=(const complex& z) -{ - real(z.real()); - imag(z.imag()); - return *this; -} -#endif - template template _CCCL_HOST_DEVICE complex& complex::operator=(const complex& z) From 8a5e56a09b88b69a760c271c8a9a6964488a71b5 Mon Sep 17 00:00:00 2001 From: Allison Piper Date: Fri, 19 Jul 2024 11:02:34 -0400 Subject: [PATCH 09/44] Add github token to nightly workflow-results action. (#2012) This is needed to fetch the runtime info from the GHA API. This only modifies the nightly workflow, the PR tests are unaffected: [skip-matrix][skip-vdc][skip-rapids][skip-docs] --- .github/workflows/ci-workflow-nightly.yml | 1 + 1 file changed, 1 insertion(+) diff --git a/.github/workflows/ci-workflow-nightly.yml b/.github/workflows/ci-workflow-nightly.yml index 619d2203d4f..fdf281b8063 100644 --- a/.github/workflows/ci-workflow-nightly.yml +++ b/.github/workflows/ci-workflow-nightly.yml @@ -133,6 +133,7 @@ jobs: id: check-workflow uses: ./.github/actions/workflow-results with: + github_token: ${{ secrets.GITHUB_TOKEN }} slack_token: ${{ secrets.SLACK_NOTIFIER_BOT_TOKEN }} slack_log: ${{ secrets.SLACK_CHANNEL_CI_LOG }} slack_alert: ${{ secrets.SLACK_CHANNEL_CI_ALERT }} From e5fcebe64471142b526427feb7c2c60e4bada562 Mon Sep 17 00:00:00 2001 From: Allison Piper Date: Sat, 20 Jul 2024 07:16:14 -0400 Subject: [PATCH 10/44] Remove obsolete build system glue from the Thrust/CUB submodule structure. (#2016) --- cub/CMakeLists.txt | 12 ++++----- cub/cmake/CubBuildTargetList.cmake | 41 ++++++++---------------------- cub/cmake/CubHeaderTesting.cmake | 4 --- cub/examples/CMakeLists.txt | 4 --- cub/test/CMakeLists.txt | 15 ----------- thrust/CMakeLists.txt | 12 +-------- 6 files changed, 16 insertions(+), 72 deletions(-) diff --git a/cub/CMakeLists.txt b/cub/CMakeLists.txt index 6e66813b7bb..91068802d63 100644 --- a/cub/CMakeLists.txt +++ b/cub/CMakeLists.txt @@ -30,7 +30,7 @@ endif() # Support adding CUB to a parent project via add_subdirectory. # See examples/cmake/add_subdir/CMakeLists.txt for details. -if (NOT CUB_TOPLEVEL_PROJECT AND NOT CUB_IN_THRUST) +if (NOT CUB_TOPLEVEL_PROJECT) include(cmake/CubAddSubdir.cmake) return() endif() @@ -51,12 +51,10 @@ mark_as_advanced(CUB_ENABLE_CPP_DIALECT_IN_NAMES) # This option is only used when CUB is built stand-alone; otherwise the Thrust # option has the same effect. -if (NOT CUB_IN_THRUST) - option(CUB_IGNORE_DEPRECATED_API - "Suppress warnings about deprecated Thrust/CUB API." - OFF - ) -endif() +option(CUB_IGNORE_DEPRECATED_API + "Suppress warnings about deprecated Thrust/CUB API." + OFF +) # Check if we're actually building anything before continuing. If not, no need # to search for deps, etc. This is a common approach for packagers that just diff --git a/cub/cmake/CubBuildTargetList.cmake b/cub/cmake/CubBuildTargetList.cmake index 426eee5e269..f01c6244f4e 100644 --- a/cub/cmake/CubBuildTargetList.cmake +++ b/cub/cmake/CubBuildTargetList.cmake @@ -132,30 +132,15 @@ function(cub_build_target_list) # Handle dialect options: set(num_dialects_enabled 0) foreach (dialect IN LISTS CUB_CPP_DIALECT_OPTIONS) - if (CUB_IN_THRUST) - # Just use Thrust's settings: - if (THRUST_ENABLE_MULTICONFIG) - set(CUB_ENABLE_DIALECT_CPP${dialect} - ${THRUST_MULTICONFIG_ENABLE_DIALECT_CPP${dialect}} - ) - else() - set(val OFF) - if (dialect EQUAL ${THRUST_CPP_DIALECT}) - set(val ON) - endif() - set(CUB_ENABLE_DIALECT_CPP${dialect} ${val}) - endif() - else() - # Create CMake options: - set(default_value OFF) - if (dialect EQUAL 14) # Default to just 14 on: - set(default_value ON) - endif() - option(CUB_ENABLE_DIALECT_CPP${dialect} - "Generate C++${dialect} build configurations." - ${default_value} - ) + # Create CMake options: + set(default_value OFF) + if (dialect EQUAL 14) # Default to just 14 on: + set(default_value ON) endif() + option(CUB_ENABLE_DIALECT_CPP${dialect} + "Generate C++${dialect} build configurations." + ${default_value} + ) if (CUB_ENABLE_DIALECT_CPP${dialect}) math(EXPR num_dialects_enabled "${num_dialects_enabled} + 1") @@ -188,14 +173,8 @@ function(cub_build_target_list) # Generic config flags: macro(add_flag_option flag docstring default) set(cub_opt "CUB_${flag}") - if (CUB_IN_THRUST) - set(thrust_opt "THRUST_${flag}") - # Use thrust's settings: - set(${cub_opt} ${${thrust_opt}}) - else() - option(${cub_opt} "${docstring}" "${default}") - mark_as_advanced(${cub_opt}) - endif() + option(${cub_opt} "${docstring}" "${default}") + mark_as_advanced(${cub_opt}) endmacro() add_flag_option(IGNORE_DEPRECATED_CPP_DIALECT "Don't warn about any deprecated C++ standards and compilers." OFF) add_flag_option(IGNORE_DEPRECATED_CPP_11 "Don't warn about deprecated C++11." OFF) diff --git a/cub/cmake/CubHeaderTesting.cmake b/cub/cmake/CubHeaderTesting.cmake index 7cead875c08..f0ca17186ce 100644 --- a/cub/cmake/CubHeaderTesting.cmake +++ b/cub/cmake/CubHeaderTesting.cmake @@ -31,10 +31,6 @@ function(cub_add_header_test label definitions) cub_clone_target_properties(${headertest_target} ${cub_target}) cub_configure_cuda_target(${headertest_target} RDC ${CUB_FORCE_RDC}) - if (CUB_IN_THRUST) - thrust_fix_clang_nvcc_build_for(${headertest_target}) - endif() - add_dependencies(cub.all.headers ${headertest_target}) add_dependencies(${config_prefix}.all ${headertest_target}) endforeach() diff --git a/cub/examples/CMakeLists.txt b/cub/examples/CMakeLists.txt index 3865b6f057f..aa766336d3e 100644 --- a/cub/examples/CMakeLists.txt +++ b/cub/examples/CMakeLists.txt @@ -35,10 +35,6 @@ function(cub_add_example target_name_var example_name example_src cub_target) cub_configure_cuda_target(${example_target} RDC ${CUB_FORCE_RDC}) target_include_directories(${example_target} PRIVATE "${CUB_SOURCE_DIR}/examples") - if (CUB_IN_THRUST) - thrust_fix_clang_nvcc_build_for(${example_target}) - endif() - # Add to the active configuration's meta target add_dependencies(${config_meta_target} ${example_target}) diff --git a/cub/test/CMakeLists.txt b/cub/test/CMakeLists.txt index 2187c66f84c..48a0142801a 100644 --- a/cub/test/CMakeLists.txt +++ b/cub/test/CMakeLists.txt @@ -202,10 +202,6 @@ function(cub_add_test target_name_var test_name test_src cub_target launcher_id) else() target_compile_definitions(${config_c2h_target} PRIVATE C2H_HAS_CURAND=0) endif() - - if (CUB_IN_THRUST) - thrust_fix_clang_nvcc_build_for(${config_c2h_target}) - endif() endif() # config_c2h_target if (CUB_SEPARATE_CATCH2) @@ -240,10 +236,6 @@ function(cub_add_test target_name_var test_name test_src cub_target launcher_id) target_link_options(${config_c2run_target} PRIVATE "-cuda") endif() - if (CUB_IN_THRUST) - thrust_fix_clang_nvcc_build_for(${config_c2run_target}) - endif() - add_test(NAME ${config_c2run_target} COMMAND "$" ) @@ -265,10 +257,6 @@ function(cub_add_test target_name_var test_name test_src cub_target launcher_id) target_compile_definitions(${test_target} PRIVATE NVRTC_CTK_PATH="-I${CUDAToolkit_INCLUDE_DIRS}") endif() - if (CUB_IN_THRUST) - thrust_fix_clang_nvcc_build_for(${test_target}) - endif() - if ("${test_target}" MATCHES "test.iterator") target_compile_options(${test_target} PRIVATE -ftemplate-depth=1000) # for handling large type lists endif() @@ -297,9 +285,6 @@ function(cub_add_test target_name_var test_name test_src cub_target launcher_id) target_link_libraries(${test_target} nvtx3-cpp) endif() - if (CUB_IN_THRUST) - thrust_fix_clang_nvcc_build_for(${test_target}) - endif() _cub_is_fail_test(is_fail_test "${test_src}") if (is_fail_test) diff --git a/thrust/CMakeLists.txt b/thrust/CMakeLists.txt index a7da7329fb6..ca372d3677f 100644 --- a/thrust/CMakeLists.txt +++ b/thrust/CMakeLists.txt @@ -58,7 +58,6 @@ option(THRUST_ENABLE_HEADER_TESTING "Test that all public headers compile." "ON" option(THRUST_ENABLE_TESTING "Build Thrust testing suite." "ON") option(THRUST_ENABLE_EXAMPLES "Build Thrust examples." "ON") option(THRUST_ENABLE_BENCHMARKS "Build Thrust runtime benchmarks." "${CCCL_ENABLE_BENCHMARKS}") -option(THRUST_INCLUDE_CUB_CMAKE "Build CUB tests and examples. (Requires CUDA)." "OFF") # Check if we're actually building anything before continuing. If not, no need # to search for deps, etc. This is a common approach for packagers that just @@ -66,8 +65,7 @@ option(THRUST_INCLUDE_CUB_CMAKE "Build CUB tests and examples. (Requires CUDA)." if (NOT (THRUST_ENABLE_HEADER_TESTING OR THRUST_ENABLE_TESTING OR THRUST_ENABLE_EXAMPLES OR - THRUST_ENABLE_BENCHMARKS OR - THRUST_INCLUDE_CUB_CMAKE)) + THRUST_ENABLE_BENCHMARKS)) return() endif() @@ -131,11 +129,3 @@ endif() if (THRUST_ENABLE_BENCHMARKS) add_subdirectory(benchmarks) endif() - -if (THRUST_INCLUDE_CUB_CMAKE AND THRUST_CUDA_FOUND) - set(CUB_IN_THRUST ON) - # CUB's path is specified generically to support both GitHub and Perforce - # source tree layouts. The include directory used by cub-config.cmake - # for source layouts is the same as the project root. - add_subdirectory("${_CUB_INCLUDE_DIR}" dependencies/cub) -endif() From 496d88d3aaa82b2af7aa1400c361cbcf14e8a118 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Mon, 22 Jul 2024 11:50:51 +0200 Subject: [PATCH 11/44] Benchmark thrust::copy with non-trivially relocatable type (#1989) --- thrust/benchmarks/bench/copy/basic.cu | 28 +++++++++++++++++++++++++-- 1 file changed, 26 insertions(+), 2 deletions(-) diff --git a/thrust/benchmarks/bench/copy/basic.cu b/thrust/benchmarks/bench/copy/basic.cu index 174a6af1178..3167c135cfb 100644 --- a/thrust/benchmarks/bench/copy/basic.cu +++ b/thrust/benchmarks/bench/copy/basic.cu @@ -37,7 +37,7 @@ static void basic(nvbench::state& state, nvbench::type_list) { const auto elements = static_cast(state.get_int64("Elements")); - thrust::device_vector input(elements, 1); + thrust::device_vector input(elements, T{1}); thrust::device_vector output(elements); state.add_element_count(elements); @@ -52,7 +52,31 @@ static void basic(nvbench::state& state, nvbench::type_list) }); } -using types = nvbench::type_list; +// Non-trivially-copyable/relocatable type which is not allowed to be copied using std::memcpy or cudaMemcpy +struct non_trivial +{ + int a; + int b; + + non_trivial() = default; + + _CCCL_HOST_DEVICE explicit non_trivial(int i) + : a(i) + , b(i) + {} + + // the user-defined copy constructor prevents the type from being trivially copyable + _CCCL_HOST_DEVICE non_trivial(const non_trivial& nt) + : a(nt.a) + , b(nt.b) + {} +}; + +static_assert(!::cuda::std::is_trivially_copyable::value, ""); // as required by the C++ standard +static_assert(!thrust::is_trivially_relocatable::value, ""); // thrust uses this check internally + +using types = + nvbench::type_list; NVBENCH_BENCH_TYPES(basic, NVBENCH_TYPE_AXES(types)) .set_name("base") From e61bafec1049b4e1b6338de1c2fba4fdf6e29a1c Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Mon, 22 Jul 2024 14:14:39 +0200 Subject: [PATCH 12/44] Make bool_constant available in C++11 (#1997) --- .../include/cuda/std/__tuple_dir/sfinae_helpers.h | 2 +- .../cuda/std/__type_traits/integral_constant.h | 11 +++-------- libcudacxx/include/cuda/std/__type_traits/is_same.h | 10 +++++----- .../include/cuda/std/__type_traits/is_signed.h | 2 +- .../include/cuda/std/__type_traits/is_unsigned.h | 2 +- .../include/cuda/std/__type_traits/negation.h | 2 +- .../include/cuda/std/detail/libcxx/include/iterator | 8 ++++---- .../include/cuda/std/detail/libcxx/include/ratio | 13 ++++++------- .../utilities/meta/meta.help/bool_constant.pass.cpp | 2 -- 9 files changed, 22 insertions(+), 30 deletions(-) diff --git a/libcudacxx/include/cuda/std/__tuple_dir/sfinae_helpers.h b/libcudacxx/include/cuda/std/__tuple_dir/sfinae_helpers.h index 88e07d83a56..f6bf304bac4 100644 --- a/libcudacxx/include/cuda/std/__tuple_dir/sfinae_helpers.h +++ b/libcudacxx/include/cuda/std/__tuple_dir/sfinae_helpers.h @@ -51,7 +51,7 @@ struct __tuple_sfinae_base {}; template - struct __test_size<__tuple_types<_Tp...>, __tuple_types<_Up...>> : _BoolConstant + struct __test_size<__tuple_types<_Tp...>, __tuple_types<_Up...>> : bool_constant {}; template