The typical invocation specifies the execution policy as follows:

thrust::inclusive_scan(thrust::device, d_input, d_input + data_size, d_output);

By passing thrust::device as the first argument, the compiler resolves this to the CUDA backend implementation defined in the library's system headers. The entry point typically looks like this:

template <typename Policy, typename InputIterator, typename OutputIterator>
OutputIterator inclusive_scan(
    const thrust::detail::execution_policy_base<Policy>& exec,
    InputIterator first,
    InputIterator last,
    OutputIterator result);

This interface forwards the request to the internal CUDA implementation found in thrust/system/cuda/detail/scan.h. Here, the implementation explicitly sets the binary operator to thrust::plus<> before delegating to a lower-level function:

template <typename Policy, typename InputIt, typename OutputIt>
OutputIt inclusive_scan(Policy& policy, InputIt first, InputIt last, OutputIt result) {
    // Default operation is addition
    return thrust::cuda_cub::inclusive_scan(policy, first, last, result, thrust::plus<>{});
}

The call chain eventually reaches the core implemantation logic, inclusive_scan_n_impl. This section reveals the critical dependency on the CUB library (CUDA Unbound). The logic determines temporary storage requirements and dispatches the kernel:

template <typename Policy, typename InputIt, typename Size, typename OutputIt, typename ScanOp>
OutputIt inclusive_scan_n_impl(Policy& policy, InputIt first, Size num_items, OutputIt result, ScanOp scan_op) {
    using ValueT = typename thrust::iterator_traits<InputIt>::value_type;
    
    // CUB Dispatch types for 32-bit and 64-bit addressing
    using Dispatch32 = cub::DispatchScan<InputIt, OutputIt, ScanOp, cub::NullType, int32_t, ValueT>;
    using Dispatch64 = cub::DispatchScan<InputIt, OutputIt, ScanOp, cub::NullType, int64_t, ValueT>;

    cudaStream_t stream = thrust::cuda_cub::stream(policy);
    cudaError_t status;

    // 1. Determine temporary storage requirements
    size_t temp_storage_bytes = 0;
    Dispatch32::Dispatch(nullptr, temp_storage_bytes, first, result, scan_op, cub::NullType{}, num_items, stream);

    // 2. Allocate temporary storage
    thrust::detail::temporary_array<std::uint8_t, Policy> temp_buffer(policy, temp_storage_bytes);

    // 3. Execute the scan operation
    Dispatch32::Dispatch(temp_buffer.data().get(), temp_storage_bytes, first, result, scan_op, cub::NullType{}, num_items, stream);

    return result + num_items;
}

As shown in the code above, Thrust delegates the actual computation to cub::DispatchScan. CUB is a header-only C++ template library that provides state-of-the-art parallel primitives. However, its scan algorithms require a temporary memory buffer, which is allocated here using thrust::detail::temporary_array. This buffer resides in global memory. The necessity to allocate and access this global memory buffer explains why Thrust's initial run might show different performance characteristics compared to a custom kernel that optimizes strictly for shared memory usage.