dylanyunlong/project_6
1
0
Fork 0
Code Issues Pull Requests Actions Projects Releases Wiki Activity

feat: CCCL CachingDeviceAllocator preload — 完整依赖链 288 files

Browse Source 
从 cccl_upstream 递归追踪 cub/util_allocator.cuh 的全部 include 依赖:
  cub/         9 files (config, util_*, version, detect_cuda_runtime)
  cuda/        libcudacxx type_traits, concepts, algorithm, iterator...
  nv/          target macros, preprocessor

总计 288 个头文件 (1.4MB)，打包到 include/ 目录，编译时 -I include
即可完全脱离 CCCL 原始目录结构。

.cu 文件直接 #include <cub/util_allocator.cuh>，
走原版 CUB CachingDeviceAllocator，零 mock。

BI-V100 参数: growth=2 bins=[8..32] max_cached=8GB/device
This commit is contained in:
dylanyunlon
2026-08-13 09:24:42 +00:00
parent 967d572073
commit 8d6f9eaeb0
290 changed files with 39904 additions and 455 deletions
Download Patch File Download Diff File Expand all files Collapse all files

29
qwen3_6_scripts/cccl_preload/include/cub/config.cuh Normal file
View File

@@ -0,0 +1,29 @@
// SPDX-FileCopyrightText: Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
// SPDX-License-Identifier: BSD-3
/**
* \file
* Static configuration header for the CUB project.
*/
#pragma once
// For _CCCL_IMPLICIT_SYSTEM_HEADER
#include <cuda/__cccl_config> // IWYU pragma: export
#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
# pragma GCC system_header
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
# pragma clang system_header
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
# pragma system_header
#endif // no system header
#include <cub/util_arch.cuh> // IWYU pragma: export
#include <cub/util_cpp_dialect.cuh> // IWYU pragma: export
#include <cub/util_macro.cuh> // IWYU pragma: export
#include <cub/util_namespace.cuh> // IWYU pragma: export
#if !_CCCL_COMPILER(NVRTC)
# include <cuda/__nvtx/nvtx.h>
#endif // !_CCCL_COMPILER(NVRTC)

45
qwen3_6_scripts/cccl_preload/include/cub/detail/detect_cuda_runtime.cuh Normal file
View File

@@ -0,0 +1,45 @@
// SPDX-FileCopyrightText: Copyright (c) 2011, Duane Merrill. All rights reserved.
// SPDX-FileCopyrightText: Copyright (c) 2011-2021, NVIDIA CORPORATION. All rights reserved.
// SPDX-License-Identifier: BSD-3
/**
* @file
* Utilities for CUDA dynamic parallelism.
*/
#pragma once
// We cannot use `cub/config.cuh` here due to circular dependencies
#include <cuda/__cccl_config>
#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
# pragma GCC system_header
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
# pragma clang system_header
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
# pragma system_header
#endif // no system header
#ifdef _CCCL_DOXYGEN_INVOKED // Only parse this during doxygen passes:
//! Defined if RDC is enabled and CUB_DISABLE_CDP is not defined.
//! Deprecated [Since 3.2]
# define CUB_RDC_ENABLED
//! If defined, support for device-side usage of CUB is disabled.
//! Deprecated [Since 3.2]. Use CCCL_DISABLE_CDP instead.
# define CUB_DISABLE_CDP
//! Execution space for functions that use the CUDA runtime API, e.g. to launch kernels. Such functions are `__host__
//! __device__` when compiling with RDC, otherwise only `__host__`.
//! Deprecated [Since 3.2]
# define CUB_RUNTIME_FUNCTION
#else // Non-doxygen pass:
# if _CCCL_HAS_CDP()
# define CUB_RDC_ENABLED
# endif // _CCCL_HAS_CDP()
# ifndef CUB_RUNTIME_FUNCTION
# define CUB_RUNTIME_FUNCTION _CCCL_CDP_API
# endif // CUB_RUNTIME_FUNCTION predefined
#endif // Do not document

901
qwen3_6_scripts/cccl_preload/include/cub/util_allocator.cuh Normal file
View File

@@ -0,0 +1,901 @@
// SPDX-FileCopyrightText: Copyright (c) 2011, Duane Merrill. All rights reserved.
// SPDX-FileCopyrightText: Copyright (c) 2011-2018, NVIDIA CORPORATION. All rights reserved.
// SPDX-License-Identifier: BSD-3
/******************************************************************************
* Simple caching allocator for device memory allocations. The allocator is
* thread-safe and capable of managing device allocations on multiple devices.
******************************************************************************/
#pragma once
#include <cub/config.cuh>
#ifndef CCCL_DISABLE_NVRTC_COMPATIBILITY_CHECK
# if _CCCL_COMPILER(NVRTC)
# error \
"Including <cub/util_allocator.cuh> is not supported when compiling with NVRTC, which supports device code only. You can define CCCL_DISABLE_NVRTC_COMPATIBILITY_CHECK to disable this warning."
# endif // _CCCL_COMPILER(NVRTC)
#endif // CCCL_DISABLE_NVRTC_COMPATIBILITY_CHECK
#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
# pragma GCC system_header
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
# pragma clang system_header
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
# pragma system_header
#endif // no system header
#include <cub/util_debug.cuh>
#include <cub/util_namespace.cuh>
#include <cuda/std/__host_stdlib/math.h>
#include <map>
#include <mutex>
#include <set>
CUB_NAMESPACE_BEGIN
/******************************************************************************
* CachingDeviceAllocator (host use)
******************************************************************************/
/**
* @brief A simple caching allocator for device memory allocations.
*
* @par Overview
* The allocator is thread-safe and stream-safe and is capable of managing cached
* device allocations on multiple devices. It behaves as follows:
*
* @par
* - Allocations from the allocator are associated with an @p active_stream. Once freed,
* the allocation becomes available immediately for reuse within the @p active_stream
* with which it was associated with during allocation, and it becomes available for
* reuse within other streams when all prior work submitted to @p active_stream has completed.
* - Allocations are categorized and cached by bin size. A new allocation request of
* a given size will only consider cached allocations within the corresponding bin.
* - Bin limits progress geometrically in accordance with the growth factor
* @p bin_growth provided during construction. Unused device allocations within
* a larger bin cache are not reused for allocation requests that categorize to
* smaller bin sizes.
* - Allocation requests below ( @p bin_growth ^ @p min_bin ) are rounded up to
* ( @p bin_growth ^ @p min_bin ).
* - Allocations above ( @p bin_growth ^ @p max_bin ) are not rounded up to the nearest
* bin and are simply freed when they are deallocated instead of being returned
* to a bin-cache.
* - If the total storage of cached allocations on a given device will exceed
* @p max_cached_bytes, allocations for that device are simply freed when they are
* deallocated instead of being returned to their bin-cache.
*
* @par
* For example, the default-constructed CachingDeviceAllocator is configured with:
* - @p bin_growth = 8
* - @p min_bin = 3
* - @p max_bin = 7
* - @p max_cached_bytes = 6MB - 1B
*
* @par
* which delineates five bin-sizes: 512B, 4KB, 32KB, 256KB, and 2MB
* and sets a maximum of 6,291,455 cached bytes per device
*
*/
struct CachingDeviceAllocator
{
//---------------------------------------------------------------------
// Constants
//---------------------------------------------------------------------
/// Out-of-bounds bin
static constexpr unsigned int INVALID_BIN = (unsigned int) -1;
/// Invalid size
static constexpr size_t INVALID_SIZE = (size_t) -1;
#ifndef _CCCL_DOXYGEN_INVOKED // Do not document
/// Invalid device ordinal
static constexpr int INVALID_DEVICE_ORDINAL = -1;
//---------------------------------------------------------------------
// Type definitions and helper types
//---------------------------------------------------------------------
/**
* Descriptor for device memory allocations
*/
struct BlockDescriptor
{
// Device pointer
void* d_ptr;
// Size of allocation in bytes
size_t bytes;
// Bin enumeration
unsigned int bin;
// device ordinal
int device;
// Associated associated_stream
cudaStream_t associated_stream;
// Signal when associated stream has run to the point at which this block was freed
cudaEvent_t ready_event;
// Constructor (suitable for searching maps for a specific block, given its pointer and
// device)
BlockDescriptor(void* d_ptr, int device)
: d_ptr(d_ptr)
, bytes(0)
, bin(INVALID_BIN)
, device(device)
, associated_stream(nullptr)
, ready_event(nullptr)
{}
// Constructor (suitable for searching maps for a range of suitable blocks, given a device)
BlockDescriptor(int device)
: d_ptr(nullptr)
, bytes(0)
, bin(INVALID_BIN)
, device(device)
, associated_stream(nullptr)
, ready_event(nullptr)
{}
// Comparison functor for comparing device pointers
static bool PtrCompare(const BlockDescriptor& a, const BlockDescriptor& b)
{
if (a.device == b.device)
{
return (a.d_ptr < b.d_ptr);
}
else
{
return (a.device < b.device);
}
}
// Comparison functor for comparing allocation sizes
static bool SizeCompare(const BlockDescriptor& a, const BlockDescriptor& b)
{
if (a.device == b.device)
{
return (a.bytes < b.bytes);
}
else
{
return (a.device < b.device);
}
}
};
/// BlockDescriptor comparator function interface
using Compare = bool (*)(const BlockDescriptor&, const BlockDescriptor&);
class TotalBytes
{
public:
size_t free;
size_t live;
TotalBytes()
{
free = live = 0;
}
};
/// Set type for cached blocks (ordered by size)
using CachedBlocks = std::multiset<BlockDescriptor, Compare>;
/// Set type for live blocks (ordered by ptr)
using BusyBlocks = std::multiset<BlockDescriptor, Compare>;
/// Map type of device ordinals to the number of cached bytes cached by each device
using GpuCachedBytes = std::map<int, TotalBytes>;
//---------------------------------------------------------------------
// Utility functions
//---------------------------------------------------------------------
/**
* Integer pow function for unsigned base and exponent
*/
static unsigned int IntPow(unsigned int base, unsigned int exp)
{
unsigned int retval = 1;
while (exp > 0)
{
if (exp & 1)
{
retval = retval * base; // multiply the result by the current base
}
base = base * base; // square the base
exp = exp >> 1; // divide the exponent in half
}
return retval;
}
/**
* Round up to the nearest power-of
*/
void NearestPowerOf(unsigned int& power, size_t& rounded_bytes, unsigned int base, size_t value)
{
power = 0;
rounded_bytes = 1;
if (value * base < value)
{
// Overflow
power = sizeof(size_t) * 8;
rounded_bytes = size_t(0) - 1;
return;
}
while (rounded_bytes < value)
{
rounded_bytes *= base;
power++;
}
}
//---------------------------------------------------------------------
// Fields
//---------------------------------------------------------------------
/// Mutex for thread-safety
std::mutex mutex;
/// Geometric growth factor for bin-sizes
unsigned int bin_growth;
/// Minimum bin enumeration
unsigned int min_bin;
/// Maximum bin enumeration
unsigned int max_bin;
/// Minimum bin size
size_t min_bin_bytes;
/// Maximum bin size
size_t max_bin_bytes;
/// Maximum aggregate cached bytes per device
size_t max_cached_bytes;
/// Whether or not to skip a call to FreeAllCached() when destructor is called.
/// (The CUDA runtime may have already shut down for statically declared allocators)
const bool skip_cleanup;
/// Whether or not to print (de)allocation events to stdout
bool debug;
/// Map of device ordinal to aggregate cached bytes on that device
GpuCachedBytes cached_bytes;
/// Set of cached device allocations available for reuse
CachedBlocks cached_blocks;
/// Set of live device allocations currently in use
BusyBlocks live_blocks;
#endif // _CCCL_DOXYGEN_INVOKED
//---------------------------------------------------------------------
// Methods
//---------------------------------------------------------------------
/**
* @brief Constructor.
*
* @param bin_growth
* Geometric growth factor for bin-sizes
*
* @param min_bin
* Minimum bin (default is bin_growth ^ 1)
*
* @param max_bin
* Maximum bin (default is no max bin)
*
* @param max_cached_bytes
* Maximum aggregate cached bytes per device (default is no limit)
*
* @param skip_cleanup
* Whether or not to skip a call to @p FreeAllCached() when the destructor is called (default
* is to deallocate)
*/
CachingDeviceAllocator(
unsigned int bin_growth,
unsigned int min_bin = 1,
unsigned int max_bin = INVALID_BIN,
size_t max_cached_bytes = INVALID_SIZE,
bool skip_cleanup = false)
: bin_growth(bin_growth)
, min_bin(min_bin)
, max_bin(max_bin)
, min_bin_bytes(IntPow(bin_growth, min_bin))
, max_bin_bytes(IntPow(bin_growth, max_bin))
, max_cached_bytes(max_cached_bytes)
, skip_cleanup(skip_cleanup)
, debug(false)
, cached_blocks(BlockDescriptor::SizeCompare)
, live_blocks(BlockDescriptor::PtrCompare)
{}
/**
* @brief Default constructor.
*
* Configured with:
* @par
* - @p bin_growth = 8
* - @p min_bin = 3
* - @p max_bin = 7
* - @p max_cached_bytes = ( @p bin_growth ^ @p max_bin) * 3 ) - 1 = 6,291,455 bytes
*
* which delineates five bin-sizes: 512B, 4KB, 32KB, 256KB, and 2MB and
* sets a maximum of 6,291,455 cached bytes per device
*/
CachingDeviceAllocator(bool skip_cleanup = false, bool debug = false)
: bin_growth(8)
, min_bin(3)
, max_bin(7)
, min_bin_bytes(IntPow(bin_growth, min_bin))
, max_bin_bytes(IntPow(bin_growth, max_bin))
, max_cached_bytes((max_bin_bytes * 3) - 1)
, skip_cleanup(skip_cleanup)
, debug(debug)
, cached_blocks(BlockDescriptor::SizeCompare)
, live_blocks(BlockDescriptor::PtrCompare)
{}
/**
* @brief Sets the limit on the number bytes this allocator is allowed to cache per device.
*
* Changing the ceiling of cached bytes does not cause any allocations (in-use or
* cached-in-reserve) to be freed. See \p FreeAllCached().
*/
cudaError_t SetMaxCachedBytes(size_t max_cached_bytes_)
{
// Lock
mutex.lock();
#ifdef CUB_DEBUG_LOG
_CubLog(
"Changing max_cached_bytes (%lld -> %lld)\n", (long long) this->max_cached_bytes, (long long) max_cached_bytes_);
#endif
this->max_cached_bytes = max_cached_bytes_;
// Unlock
mutex.unlock();
return cudaSuccess;
}
/**
* @brief Provides a suitable allocation of device memory for the given size on the specified
* device.
*
* Once freed, the allocation becomes available immediately for reuse within the @p
* active_stream with which it was associated with during allocation, and it becomes available
* for reuse within other streams when all prior work submitted to @p active_stream has
* completed.
*
* @param[in] device
* Device on which to place the allocation
*
* @param[out] d_ptr
* Reference to pointer to the allocation
*
* @param[in] bytes
* Minimum number of bytes for the allocation
*
* @param[in] active_stream
* The stream to be associated with this allocation
*/
cudaError_t DeviceAllocate(int device, void** d_ptr, size_t bytes, cudaStream_t active_stream = nullptr)
{
*d_ptr = nullptr;
int entrypoint_device = INVALID_DEVICE_ORDINAL;
cudaError_t error = cudaSuccess;
if (device == INVALID_DEVICE_ORDINAL)
{
error = CubDebug(cudaGetDevice(&entrypoint_device));
if (cudaSuccess != error)
{
return error;
}
device = entrypoint_device;
}
// Create a block descriptor for the requested allocation
bool found = false;
BlockDescriptor search_key(device);
search_key.associated_stream = active_stream;
NearestPowerOf(search_key.bin, search_key.bytes, bin_growth, bytes);
if (search_key.bin > max_bin)
{
// Bin is greater than our maximum bin: allocate the request
// exactly and give out-of-bounds bin. It will not be cached
// for reuse when returned.
search_key.bin = INVALID_BIN;
search_key.bytes = bytes;
}
else
{
// Search for a suitable cached allocation: lock
mutex.lock();
if (search_key.bin < min_bin)
{
// Bin is less than minimum bin: round up
search_key.bin = min_bin;
search_key.bytes = min_bin_bytes;
}
// Iterate through the range of cached blocks on the same device in the same bin
CachedBlocks::iterator block_itr = cached_blocks.lower_bound(search_key);
while ((block_itr != cached_blocks.end()) && (block_itr->device == device) && (block_itr->bin == search_key.bin))
{
// To prevent races with reusing blocks returned by the host but still
// in use by the device, only consider cached blocks that are
// either (from the active stream) or (from an idle stream)
bool is_reusable = false;
if (active_stream == block_itr->associated_stream)
{
is_reusable = true;
}
else
{
const cudaError_t event_status = cudaEventQuery(block_itr->ready_event);
if (event_status != cudaErrorNotReady)
{
CubDebug(event_status);
is_reusable = true;
}
}
if (is_reusable)
{
// Reuse existing cache block. Insert into live blocks.
found = true;
search_key = *block_itr;
search_key.associated_stream = active_stream;
live_blocks.insert(search_key);
// Remove from free blocks
cached_bytes[device].free -= search_key.bytes;
cached_bytes[device].live += search_key.bytes;
#ifdef CUB_DEBUG_LOG
_CubLog("\tDevice %d reused cached block at %p (%lld bytes) for stream %lld (previously associated with "
"stream %lld).\n",
device,
search_key.d_ptr,
(long long) search_key.bytes,
(long long) search_key.associated_stream,
(long long) block_itr->associated_stream);
#endif
cached_blocks.erase(block_itr);
break;
}
block_itr++;
}
// Done searching: unlock
mutex.unlock();
}
// Allocate the block if necessary
if (!found)
{
// Set runtime's current device to specified device (entrypoint may not be set)
if (device != entrypoint_device)
{
error = CubDebug(cudaGetDevice(&entrypoint_device));
if (cudaSuccess != error)
{
return error;
}
error = CubDebug(cudaSetDevice(device));
if (cudaSuccess != error)
{
return error;
}
}
// Attempt to allocate
error = CubDebug(cudaMalloc(&search_key.d_ptr, search_key.bytes));
if (error == cudaErrorMemoryAllocation)
{
// The allocation attempt failed: free all cached blocks on device and retry
#ifdef CUB_DEBUG_LOG
_CubLog("\tDevice %d failed to allocate %lld bytes for stream %lld, retrying after freeing cached allocations",
device,
(long long) search_key.bytes,
(long long) search_key.associated_stream);
#endif
error = cudaSuccess; // Reset the error we will return
cudaGetLastError(); // Reset CUDART's error
// Lock
mutex.lock();
// Iterate the range of free blocks on the same device
BlockDescriptor free_key(device);
CachedBlocks::iterator block_itr = cached_blocks.lower_bound(free_key);
while ((block_itr != cached_blocks.end()) && (block_itr->device == device))
{
// No need to worry about synchronization with the device: cudaFree is
// blocking and will synchronize across all kernels executing
// on the current device
// Free device memory and destroy stream event.
error = CubDebug(cudaFree(block_itr->d_ptr));
if (cudaSuccess != error)
{
break;
}
error = CubDebug(cudaEventDestroy(block_itr->ready_event));
if (cudaSuccess != error)
{
break;
}
// Reduce balance and erase entry
cached_bytes[device].free -= block_itr->bytes;
#ifdef CUB_DEBUG_LOG
_CubLog("\tDevice %d freed %lld bytes.\n\t\t %lld available blocks cached (%lld bytes), %lld live blocks "
"(%lld bytes) outstanding.\n",
device,
(long long) block_itr->bytes,
(long long) cached_blocks.size(),
(long long) cached_bytes[device].free,
(long long) live_blocks.size(),
(long long) cached_bytes[device].live);
#endif
block_itr = cached_blocks.erase(block_itr);
}
// Unlock
mutex.unlock();
// Return under error
if (error)
{
return error;
}
// Try to allocate again
error = CubDebug(cudaMalloc(&search_key.d_ptr, search_key.bytes));
if (cudaSuccess != error)
{
return error;
}
}
// Create ready event
error = CubDebug(cudaEventCreateWithFlags(&search_key.ready_event, cudaEventDisableTiming));
if (cudaSuccess != error)
{
return error;
}
// Insert into live blocks
mutex.lock();
live_blocks.insert(search_key);
cached_bytes[device].live += search_key.bytes;
mutex.unlock();
#ifdef CUB_DEBUG_LOG
_CubLog("\tDevice %d allocated new device block at %p (%lld bytes associated with stream %lld).\n",
device,
search_key.d_ptr,
(long long) search_key.bytes,
(long long) search_key.associated_stream);
#endif
// Attempt to revert back to previous device if necessary
if ((entrypoint_device != INVALID_DEVICE_ORDINAL) && (entrypoint_device != device))
{
error = CubDebug(cudaSetDevice(entrypoint_device));
if (cudaSuccess != error)
{
return error;
}
}
}
// Copy device pointer to output parameter
*d_ptr = search_key.d_ptr;
#ifdef CUB_DEBUG_LOG
if (debug)
{
_CubLog("\t\t%lld available blocks cached (%lld bytes), %lld live blocks outstanding(%lld bytes).\n",
(long long) cached_blocks.size(),
(long long) cached_bytes[device].free,
(long long) live_blocks.size(),
(long long) cached_bytes[device].live);
}
#endif
return error;
}
/**
* @brief Provides a suitable allocation of device memory for the given size on the current
* device.
*
* Once freed, the allocation becomes available immediately for reuse within the @p
* active_stream with which it was associated with during allocation, and it becomes available
* for reuse within other streams when all prior work submitted to @p active_stream has
* completed.
*
* @param[out] d_ptr
* Reference to pointer to the allocation
*
* @param[in] bytes
* Minimum number of bytes for the allocation
*
* @param[in] active_stream
* The stream to be associated with this allocation
*/
cudaError_t DeviceAllocate(void** d_ptr, size_t bytes, cudaStream_t active_stream = nullptr)
{
return DeviceAllocate(INVALID_DEVICE_ORDINAL, d_ptr, bytes, active_stream);
}
/**
* @brief Frees a live allocation of device memory on the specified device, returning it to the
* allocator.
*
* Once freed, the allocation becomes available immediately for reuse within the
* @p active_stream with which it was associated with during allocation, and it becomes
* available for reuse within other streams when all prior work submitted to @p active_stream
* has completed.
*/
cudaError_t DeviceFree(int device, void* d_ptr)
{
int entrypoint_device = INVALID_DEVICE_ORDINAL;
cudaError_t error = cudaSuccess;
if (device == INVALID_DEVICE_ORDINAL)
{
error = CubDebug(cudaGetDevice(&entrypoint_device));
if (cudaSuccess != error)
{
return error;
}
device = entrypoint_device;
}
// Lock
mutex.lock();
// Find corresponding block descriptor
bool recached = false;
BlockDescriptor search_key(d_ptr, device);
BusyBlocks::iterator block_itr = live_blocks.find(search_key);
if (block_itr != live_blocks.end())
{
// Remove from live blocks
search_key = *block_itr;
live_blocks.erase(block_itr);
cached_bytes[device].live -= search_key.bytes;
// Keep the returned allocation if bin is valid and we won't exceed the max cached threshold
if ((search_key.bin != INVALID_BIN) && (cached_bytes[device].free + search_key.bytes <= max_cached_bytes))
{
// Insert returned allocation into free blocks
recached = true;
cached_blocks.insert(search_key);
cached_bytes[device].free += search_key.bytes;
#ifdef CUB_DEBUG_LOG
_CubLog("\tDevice %d returned %lld bytes from associated stream %lld.\n\t\t %lld available blocks cached (%lld "
"bytes), %lld live blocks outstanding. (%lld bytes)\n",
device,
(long long) search_key.bytes,
(long long) search_key.associated_stream,
(long long) cached_blocks.size(),
(long long) cached_bytes[device].free,
(long long) live_blocks.size(),
(long long) cached_bytes[device].live);
#endif
}
}
// Unlock
mutex.unlock();
// First set to specified device (entrypoint may not be set)
if (device != entrypoint_device)
{
error = CubDebug(cudaGetDevice(&entrypoint_device));
if (cudaSuccess != error)
{
return error;
}
error = CubDebug(cudaSetDevice(device));
if (cudaSuccess != error)
{
return error;
}
}
if (recached)
{
// Insert the ready event in the associated stream (must have current device set properly)
error = CubDebug(cudaEventRecord(search_key.ready_event, search_key.associated_stream));
if (cudaSuccess != error)
{
return error;
}
}
if (!recached)
{
// Free the allocation from the runtime and cleanup the event.
error = CubDebug(cudaFree(d_ptr));
if (cudaSuccess != error)
{
return error;
}
error = CubDebug(cudaEventDestroy(search_key.ready_event));
if (cudaSuccess != error)
{
return error;
}
#ifdef CUB_DEBUG_LOG
_CubLog("\tDevice %d freed %lld bytes from associated stream %lld.\n\t\t %lld available blocks cached (%lld "
"bytes), %lld live blocks (%lld bytes) outstanding.\n",
device,
(long long) search_key.bytes,
(long long) search_key.associated_stream,
(long long) cached_blocks.size(),
(long long) cached_bytes[device].free,
(long long) live_blocks.size(),
(long long) cached_bytes[device].live);
#endif
}
// Reset device
if ((entrypoint_device != INVALID_DEVICE_ORDINAL) && (entrypoint_device != device))
{
error = CubDebug(cudaSetDevice(entrypoint_device));
if (cudaSuccess != error)
{
return error;
}
}
return error;
}
/**
* @brief Frees a live allocation of device memory on the current device, returning it to the
* allocator.
*
* Once freed, the allocation becomes available immediately for reuse within the @p
* active_stream with which it was associated with during allocation, and it becomes available
* for reuse within other streams when all prior work submitted to @p active_stream has
* completed.
*/
cudaError_t DeviceFree(void* d_ptr)
{
return DeviceFree(INVALID_DEVICE_ORDINAL, d_ptr);
}
/**
* @brief Frees all cached device allocations on all devices
*/
cudaError_t FreeAllCached()
{
cudaError_t error = cudaSuccess;
int entrypoint_device = INVALID_DEVICE_ORDINAL;
int current_device = INVALID_DEVICE_ORDINAL;
mutex.lock();
while (!cached_blocks.empty())
{
// Get first block
CachedBlocks::iterator begin = cached_blocks.begin();
// Get entry-point device ordinal if necessary
if (entrypoint_device == INVALID_DEVICE_ORDINAL)
{
error = CubDebug(cudaGetDevice(&entrypoint_device));
if (cudaSuccess != error)
{
break;
}
}
// Set current device ordinal if necessary
if (begin->device != current_device)
{
error = CubDebug(cudaSetDevice(begin->device));
if (cudaSuccess != error)
{
break;
}
current_device = begin->device;
}
// Free device memory
error = CubDebug(cudaFree(begin->d_ptr));
if (cudaSuccess != error)
{
break;
}
error = CubDebug(cudaEventDestroy(begin->ready_event));
if (cudaSuccess != error)
{
break;
}
// Reduce balance and erase entry
const size_t block_bytes = begin->bytes;
cached_bytes[current_device].free -= block_bytes;
cached_blocks.erase(begin);
#ifdef CUB_DEBUG_LOG
_CubLog("\tDevice %d freed %lld bytes.\n\t\t %lld available blocks cached (%lld bytes), %lld live blocks (%lld "
"bytes) outstanding.\n",
current_device,
(long long) block_bytes,
(long long) cached_blocks.size(),
(long long) cached_bytes[current_device].free,
(long long) live_blocks.size(),
(long long) cached_bytes[current_device].live);
#endif
}
mutex.unlock();
// Attempt to revert back to entry-point device if necessary
if (entrypoint_device != INVALID_DEVICE_ORDINAL)
{
error = CubDebug(cudaSetDevice(entrypoint_device));
if (cudaSuccess != error)
{
return error;
}
}
return error;
}
/**
* @brief Destructor
*/
virtual ~CachingDeviceAllocator()
{
if (!skip_cleanup)
{
FreeAllCached();
}
}
};
CUB_NAMESPACE_END

219
qwen3_6_scripts/cccl_preload/include/cub/util_arch.cuh Normal file
View File

@@ -0,0 +1,219 @@
// SPDX-FileCopyrightText: Copyright (c) 2011, Duane Merrill. All rights reserved.
// SPDX-FileCopyrightText: Copyright (c) 2011-2025, NVIDIA CORPORATION. All rights reserved.
// SPDX-License-Identifier: BSD-3
/**
* \file
* Static architectural properties by SM version.
*/
#pragma once
#include <cub/config.cuh>
#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
# pragma GCC system_header
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
# pragma clang system_header
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
# pragma system_header
#endif // no system header
#include <cub/util_cpp_dialect.cuh> // IWYU pragma: export
#include <cub/util_macro.cuh>
#include <cub/util_namespace.cuh>
#include <cuda/__cmath/ceil_div.h>
#include <cuda/__cmath/round_up.h>
#include <cuda/__device/compute_capability.h>
#include <cuda/std/__algorithm/clamp.h>
#include <cuda/std/__algorithm/max.h>
#include <cuda/std/__algorithm/min.h>
// Legacy include; this functionality used to be defined in here.
#include <cub/detail/detect_cuda_runtime.cuh>
CUB_NAMESPACE_BEGIN
#ifndef _CCCL_DOXYGEN_INVOKED // Do not document
/// In device code, CUB_PTX_ARCH expands to the PTX version for which we are
/// compiling. In host code, CUB_PTX_ARCH's value is implementation defined.
# ifndef CUB_PTX_ARCH
// deprecated in 3.1
# if _CCCL_CUDA_COMPILER(NVHPC)
// NV_TARGET_MINIMUM_SM_INTEGER is the oldest target PTX version, and is defined when compiling both host code and
// device code.
# define CUB_PTX_ARCH (NV_TARGET_MINIMUM_SM_INTEGER * 10)
# else // ^^^ _CCCL_CUDA_COMPILER(NVHPC) ^^^ / vvv !_CCCL_CUDA_COMPILER(NVHPC) vvv
# define CUB_PTX_ARCH _CCCL_PTX_ARCH()
# endif // ^^^ !_CCCL_CUDA_COMPILER(NVHPC) ^^^
# endif
/// Maximum number of devices supported.
# ifndef CUB_MAX_DEVICES
//! Deprecated [Since 3.0]
# define CUB_MAX_DEVICES (128)
# endif
static_assert(CUB_MAX_DEVICES > 0, "CUB_MAX_DEVICES must be greater than 0.");
/// Number of threads per warp
# ifndef CUB_LOG_WARP_THREADS
//! Deprecated [Since 3.0]
# define CUB_LOG_WARP_THREADS(unused) (5)
//! Deprecated [Since 3.0]
# define CUB_WARP_THREADS(unused) (1 << CUB_LOG_WARP_THREADS(0))
//! Deprecated [Since 3.0]
# define CUB_PTX_WARP_THREADS CUB_WARP_THREADS(0)
//! Deprecated [Since 3.0]
# define CUB_PTX_LOG_WARP_THREADS CUB_LOG_WARP_THREADS(0)
# endif
/// Number of smem banks
# ifndef CUB_LOG_SMEM_BANKS
//! Deprecated [Since 3.0]
# define CUB_LOG_SMEM_BANKS(unused) (5)
//! Deprecated [Since 3.0]
# define CUB_SMEM_BANKS(unused) (1 << CUB_LOG_SMEM_BANKS(0))
//! Deprecated [Since 3.0]
# define CUB_PTX_LOG_SMEM_BANKS CUB_LOG_SMEM_BANKS(0)
//! Deprecated [Since 3.0]
# define CUB_PTX_SMEM_BANKS CUB_SMEM_BANKS
# endif
/// Oversubscription factor
# ifndef CUB_SUBSCRIPTION_FACTOR
//! Deprecated [Since 3.0]
# define CUB_SUBSCRIPTION_FACTOR(unused) (5)
//! Deprecated [Since 3.0]
# define CUB_PTX_SUBSCRIPTION_FACTOR CUB_SUBSCRIPTION_FACTOR(0)
# endif
/// Prefer padding overhead vs X-way conflicts greater than this threshold
# ifndef CUB_PREFER_CONFLICT_OVER_PADDING
//! Deprecated [Since 3.0]
# define CUB_PREFER_CONFLICT_OVER_PADDING(unused) (1)
//! Deprecated [Since 3.0]
# define CUB_PTX_PREFER_CONFLICT_OVER_PADDING CUB_PREFER_CONFLICT_OVER_PADDING(0)
# endif
namespace detail
{
inline constexpr int max_devices = CUB_MAX_DEVICES;
inline constexpr int warp_threads = CUB_PTX_WARP_THREADS;
inline constexpr int log2_warp_threads = CUB_PTX_LOG_WARP_THREADS;
inline constexpr int smem_banks = CUB_SMEM_BANKS(0);
inline constexpr int log2_smem_banks = CUB_PTX_LOG_SMEM_BANKS;
inline constexpr int subscription_factor = CUB_PTX_SUBSCRIPTION_FACTOR;
inline constexpr bool prefer_conflict_over_padding = CUB_PTX_PREFER_CONFLICT_OVER_PADDING;
// The maximum amount of shared memory available per thread block for eternity. Every current and future CUDA
// architecture has and will have at least this amount of shared memory. This is also the maximum size of total static
// shared memory in a kernel. Note that dynamic shared memory may be larger than this amount.
static constexpr ::cuda::std::size_t max_smem_per_block = 48 * 1024;
// The size in bytes of the largest machine word that can be atomically read/written in a single instruction, so we can
// use it to pass messages from one thread to another using strong loads (acquire) and stores (release).
inline constexpr int largest_atomic_message_size = 16;
struct scaling_result
{
int items_per_thread;
int threads_per_block;
};
[[nodiscard]] _CCCL_HOST_DEVICE_API inline constexpr auto
scale_reg_bound(int nominal_4B_threads_per_block, int nominal_4B_items_per_thread, int target_type_size)
-> scaling_result
{
const int items_per_thread =
(::cuda::std::max) (1, nominal_4B_items_per_thread * 4 / (::cuda::std::max) (4, target_type_size));
const int threads_per_block =
(::cuda::std::min) (nominal_4B_threads_per_block,
::cuda::ceil_div(int{max_smem_per_block} / (target_type_size * items_per_thread), 32) * 32);
return {items_per_thread, threads_per_block};
}
template <int Nominal4ByteThreadsPerBlock, int Nominal4ByteItemsPerThread, typename T>
struct RegBoundScaling
{
private:
static constexpr auto result =
scale_reg_bound(Nominal4ByteThreadsPerBlock, Nominal4ByteItemsPerThread, int{sizeof(T)});
public:
static constexpr int ITEMS_PER_THREAD = result.items_per_thread;
static constexpr int BLOCK_THREADS = result.threads_per_block;
};
[[nodiscard]] _CCCL_HOST_DEVICE_API inline constexpr auto
scale_mem_bound(int nominal_4B_threads_per_block, int nominal_4B_items_per_thread, int target_type_size)
-> scaling_result
{
const int items_per_thread =
::cuda::std::clamp(nominal_4B_items_per_thread * 4 / target_type_size, 1, nominal_4B_items_per_thread * 2);
const int threads_per_block =
(::cuda::std::min) (nominal_4B_threads_per_block,
::cuda::round_up(int{max_smem_per_block} / (target_type_size * items_per_thread), 32));
return {items_per_thread, threads_per_block};
}
template <int Nominal4ByteThreadsPerBlock, int Nominal4ByteItemsPerThread, typename T>
struct MemBoundScaling
{
private:
static constexpr auto result =
scale_mem_bound(Nominal4ByteThreadsPerBlock, Nominal4ByteItemsPerThread, int{sizeof(T)});
public:
static constexpr int ITEMS_PER_THREAD = result.items_per_thread;
static constexpr int BLOCK_THREADS = result.threads_per_block;
};
template <int Nominal4ByteThreadsPerBlock, int Nominal4ByteItemsPerThread, typename = void>
struct NoScaling
{
static constexpr int ITEMS_PER_THREAD = Nominal4ByteItemsPerThread;
static constexpr int BLOCK_THREADS = Nominal4ByteThreadsPerBlock;
};
[[nodiscard]] _CCCL_HOST_DEVICE_API constexpr ::cuda::compute_capability current_tuning_cc() noexcept
{
# if _CCCL_CUDA_COMPILER(NVHPC)
return ::cuda::compute_capability(NV_TARGET_MINIMUM_SM_INTEGER);
# elif _CCCL_DEVICE_COMPILATION()
return ::cuda::device::current_compute_capability();
# else
// clang 22+ supports __CUDA_ARCH_LIST__ and also instantiates tuning policies inside kernels during the **host**
// pass (e.g. to compute the value for __launch_bounds__), where we rely on current_tuning_cc(), which is then passed
// to the policy selector. In the rare case that the policy selector is an adapter over a policy hub and invokes
// ChainedPolicy (e.g. test cub.test.device.histogram_custom_policy_hub.lid_0), it will fail to compile during
// constant evaluation, since it cannot find a policy for a PTX version of zero. As a workaround, we return the oldest
// CC we are compiling for during the host pass. And for consistency, we do the same for all compilers.
# if _CCCL_CUDA_COMPILER(CLANG)
return ::cuda::__target_compute_capabilities().front();
# else // ^^^ _CCCL_CUDA_COMPILER(CLANG) ^^^ / vvv !_CCCL_CUDA_COMPILER(CLANG) vvv
return {};
# endif // ^^^ !_CCCL_CUDA_COMPILER(CLANG) ^^^
# endif
}
_CCCL_EXEC_CHECK_DISABLE
template <class PolicySelector>
[[nodiscard]] _CCCL_HOST_DEVICE_API constexpr auto select_policy(::cuda::compute_capability cc)
{
return PolicySelector{}(cc);
}
template <class PolicySelector>
[[nodiscard]] _CCCL_DEVICE_API constexpr auto current_policy()
{
return select_policy<PolicySelector>(current_tuning_cc());
}
} // namespace detail
#endif // Do not document
CUB_NAMESPACE_END

71
qwen3_6_scripts/cccl_preload/include/cub/util_cpp_dialect.cuh Normal file
View File

@@ -0,0 +1,71 @@
// SPDX-FileCopyrightText: Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
// SPDX-License-Identifier: BSD-3
//! @file
//! Detect the version of the C++ standard used by the compiler.
#pragma once
#include <cuda/__cccl_config>
#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
# pragma GCC system_header
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
# pragma clang system_header
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
# pragma system_header
#endif // no system header
#ifndef _CCCL_DOXYGEN_INVOKED // Do not document
// Deprecation warnings may be silenced by defining the following macros. These
// may be combined.
// - CCCL_IGNORE_DEPRECATED_COMPILER
// Ignore deprecation warnings when using deprecated compilers. Compiling
// with deprecated C++ dialects will still issue warnings.
//! Deprecated [Since 3.0]
# define CUB_CPP_DIALECT _CCCL_STD_VER
// Define CUB_COMPILER_DEPRECATION macro:
# if _CCCL_COMPILER(MSVC) || _CCCL_COMPILER(NVRTC)
# define CUB_COMP_DEPR_IMPL(msg) _CCCL_PRAGMA(message(__FILE__ ":" _CCCL_TO_STRING(__LINE__) ": warning: " #msg))
# else // clang / gcc:
# define CUB_COMP_DEPR_IMPL(msg) _CCCL_PRAGMA(GCC warning #msg)
# endif
// Compiler checks:
// clang-format off
# define CUB_COMPILER_DEPRECATION(REQ) \
CUB_COMP_DEPR_IMPL(CUB requires at least REQ. Define CCCL_IGNORE_DEPRECATED_COMPILER to suppress this message.)
# define CUB_COMPILER_DEPRECATION_SOFT(REQ, CUR) \
CUB_COMP_DEPR_IMPL( \
CUB requires at least REQ. CUR is deprecated but still supported. CUR support will be removed in a \
future release. Define CCCL_IGNORE_DEPRECATED_COMPILER to suppress this message.)
// clang-format on
# ifndef CCCL_IGNORE_DEPRECATED_COMPILER
# if _CCCL_COMPILER(GCC, <, 7)
CUB_COMPILER_DEPRECATION(GCC 7.0);
# elif _CCCL_COMPILER(CLANG, <, 7)
CUB_COMPILER_DEPRECATION(Clang 7.0);
# elif _CCCL_COMPILER(MSVC, <, 19, 10)
// <2017. Hard upgrade message:
CUB_COMPILER_DEPRECATION(MSVC 2019(19.20 / 16.0 / 14.20));
# endif
# endif // CCCL_IGNORE_DEPRECATED_COMPILER
# undef CUB_COMPILER_DEPRECATION_SOFT
# undef CUB_COMPILER_DEPRECATION
// C++17 dialect check:
# ifndef CCCL_IGNORE_DEPRECATED_CPP_DIALECT
# if _CCCL_STD_VER < 2017
# error CUB requires at least C++17. Define CCCL_IGNORE_DEPRECATED_CPP_DIALECT to suppress this message.
# endif // _CCCL_STD_VER < 2017
# endif
# undef CUB_COMP_DEPR_IMPL
#endif // !_CCCL_DOXYGEN_INVOKED

187
qwen3_6_scripts/cccl_preload/include/cub/util_debug.cuh Normal file
View File

@@ -0,0 +1,187 @@
// SPDX-FileCopyrightText: Copyright (c) 2011, Duane Merrill. All rights reserved.
// SPDX-FileCopyrightText: Copyright (c) 2011-2022, NVIDIA CORPORATION. All rights reserved.
// SPDX-License-Identifier: BSD-3
/**
* \file
* Error and event logging routines.
*
* The following macros definitions are supported:
* - \p CUB_LOG. Simple event messages are printed to \p stdout.
*/
#pragma once
#include <cub/config.cuh>
#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
# pragma GCC system_header
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
# pragma clang system_header
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
# pragma system_header
#endif // no system header
#include <nv/target>
#ifdef _CCCL_DOXYGEN_INVOKED // Only parse this during doxygen passes:
/**
* @def CUB_DEBUG_LOG
*
* Causes kernel launch configurations to be printed to the console
*/
# define CUB_DEBUG_LOG
/**
* @def CUB_DEBUG_SYNC
*
* Causes synchronization of the stream after every kernel launch to check
* for errors. Also causes kernel launch configurations to be printed to the
* console.
*/
# define CUB_DEBUG_SYNC
/**
* @def CUB_DEBUG_ALL
*
* Causes host and device-side precondition assertions to be checked. Apart
* from that, causes synchronization of the stream after every kernel launch to
* check for errors. Also causes kernel launch configurations to be printed to
* the console.
*/
# define CUB_DEBUG_ALL
#endif // _CCCL_DOXYGEN_INVOKED
// CUB_DEBUG_SYNC also enables CUB_DEBUG_LOG
#ifdef CUB_DEBUG_SYNC
# ifndef CUB_DEBUG_LOG
# define CUB_DEBUG_LOG
# endif
#endif
// CUB_DEBUG_ALL = CUB_DEBUG_LOG + CUB_DEBUG_SYNC
#ifdef CUB_DEBUG_ALL
# ifndef CUB_DEBUG_LOG
# define CUB_DEBUG_LOG
# endif // CUB_DEBUG_LOG
# ifndef CUB_DEBUG_SYNC
# define CUB_DEBUG_SYNC
# endif // CUB_DEBUG_SYNC
#endif // CUB_DEBUG_ALL
/// CUB error reporting macro (prints error messages to stderr)
#if (defined(DEBUG) || defined(_DEBUG)) && !defined(CUB_STDERR)
# define CUB_STDERR
#endif
#if defined(CUB_STDERR) || defined(CUB_DEBUG_LOG)
# include <cuda/std/__host_stdlib/cstdio>
#endif
CUB_NAMESPACE_BEGIN
/**
* \brief %If \p CUB_STDERR is defined and \p error is not \p cudaSuccess, the
* corresponding error message is printed to \p stderr (or \p stdout in device
* code) along with the supplied source context.
*
* \return The CUDA error.
*/
_CCCL_HOST_DEVICE _CCCL_FORCEINLINE cudaError_t
Debug(cudaError_t error, [[maybe_unused]] const char* filename, [[maybe_unused]] int line)
{
// Clear the global CUDA error state which may have been set by the last
// call. Otherwise, errors may "leak" to unrelated kernel launches.
// clang-format off
#ifndef CUB_RDC_ENABLED
#define CUB_TEMP_DEVICE_CODE
#else
#define CUB_TEMP_DEVICE_CODE last_error = cudaGetLastError()
#endif
cudaError_t last_error = cudaSuccess;
NV_IF_ELSE_TARGET(
NV_IS_HOST,
(last_error = cudaGetLastError();),
(CUB_TEMP_DEVICE_CODE;)
);
#undef CUB_TEMP_DEVICE_CODE
// clang-format on
if (error == cudaSuccess && last_error != cudaSuccess)
{
error = last_error;
}
#ifdef CUB_STDERR
if (error)
{
NV_IF_ELSE_TARGET(
NV_IS_HOST,
(fprintf(stderr, "CUDA error %d [%s, %d]: %s\n", error, filename, line, cudaGetErrorString(error));
fflush(stderr);),
(printf("CUDA error %d [block (%d,%d,%d) thread (%d,%d,%d), %s, %d]\n",
error,
blockIdx.z,
blockIdx.y,
blockIdx.x,
threadIdx.z,
threadIdx.y,
threadIdx.x,
filename,
line);));
}
#endif
return error;
}
/**
* \brief Debug macro
*/
#ifndef CubDebug
# define CubDebug(e) CUB_NS_QUALIFIER::Debug((cudaError_t) (e), __FILE__, __LINE__)
#endif
/**
* \brief Debug macro with exit
*/
#ifndef CubDebugExit
# define CubDebugExit(e) \
if (CUB_NS_QUALIFIER::Debug((cudaError_t) (e), __FILE__, __LINE__)) \
{ \
exit(1); \
}
#endif
/**
* \brief Log macro for printf statements.
*/
#if !defined(_CubLog)
# if _CCCL_HOSTJIT()
# define _CubLog(format, ...) (void(0))
# else // ^^^ _CCCL_HOSTJIT() ^^^ / vvv !_CCCL_HOSTJIT() vvv
# define _CubLog(format, ...) \
do \
{ \
NV_IF_ELSE_TARGET( \
NV_IS_HOST, \
(printf(format, __VA_ARGS__);), \
(printf("[block (%d,%d,%d), thread (%d,%d,%d)]: " format, \
blockIdx.z, \
blockIdx.y, \
blockIdx.x, \
threadIdx.z, \
threadIdx.y, \
threadIdx.x, \
__VA_ARGS__);)); \
} while (false)
# endif // !_CCCL_HOSTJIT()
#endif // !defined(_CubLog)
CUB_NAMESPACE_END

67
qwen3_6_scripts/cccl_preload/include/cub/util_macro.cuh Normal file
View File

@@ -0,0 +1,67 @@
// SPDX-FileCopyrightText: Copyright (c) 2011, Duane Merrill. All rights reserved.
// SPDX-FileCopyrightText: Copyright (c) 2011-2024, NVIDIA CORPORATION. All rights reserved.
// SPDX-License-Identifier: BSD-3
/******************************************************************************
* Common C/C++ macro utilities
******************************************************************************/
#pragma once
#include <cuda/__cccl_config>
#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
# pragma GCC system_header
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
# pragma clang system_header
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
# pragma system_header
#endif // no system header
#include <cub/detail/detect_cuda_runtime.cuh> // IWYU pragma: export
#include <cub/util_namespace.cuh> // IWYU pragma: export
CUB_NAMESPACE_BEGIN
#ifdef _CCCL_DOXYGEN_INVOKED
# define CUB_DISABLE_KERNEL_VISIBILITY_WARNING_SUPPRESSION
#endif
/**
* @def CUB_DISABLE_KERNEL_VISIBILITY_WARNING_SUPPRESSION
* If defined, the default suppression of kernel visibility attribute warning is disabled.
*/
#if !defined(CUB_DISABLE_KERNEL_VISIBILITY_WARNING_SUPPRESSION)
_CCCL_DIAG_SUPPRESS_GCC("-Wattributes")
_CCCL_DIAG_SUPPRESS_CLANG("-Wattributes")
# if !_CCCL_CUDA_COMPILER(NVHPC)
_CCCL_DIAG_SUPPRESS_NVHPC(attribute_requires_external_linkage)
# endif // !_CCCL_CUDA_COMPILER(NVHPC)
#endif // !CUB_DISABLE_KERNEL_VISIBILITY_WARNING_SUPPRESSION
#ifndef CUB_DEFINE_KERNEL_GETTER
# define CUB_DEFINE_KERNEL_GETTER(name, ...) \
_CCCL_HIDE_FROM_ABI CUB_RUNTIME_FUNCTION static constexpr decltype(&__VA_ARGS__) name() \
{ \
return &__VA_ARGS__; \
}
#endif
// TODO(bgruber): drop in CCCL 4.0 when we drop the public dispatchers
#ifndef CUB_DEFINE_SUB_POLICY_GETTER
# define CUB_DEFINE_SUB_POLICY_GETTER(name) \
_CCCL_HOST_DEVICE static constexpr auto name() \
{ \
return MakePolicyWrapper(typename StaticPolicyT::name##Policy()); \
}
#endif
#if defined(CUB_DEFINE_RUNTIME_POLICIES)
# define CUB_DETAIL_STATIC_ISH_ASSERT(expr, msg) _CCCL_ASSERT(expr, msg)
# define CUB_DETAIL_CONSTEXPR_ISH
#else // ^^^ CUB_DEFINE_RUNTIME_POLICIES ^^^ / vvv !CUB_DEFINE_RUNTIME_POLICIES vvv
# define CUB_DETAIL_STATIC_ISH_ASSERT(expr, msg) static_assert(expr, msg);
# define CUB_DETAIL_CONSTEXPR_ISH constexpr
#endif // !(CUB_DEFINE_RUNTIME_POLICIES)
CUB_NAMESPACE_END

172
qwen3_6_scripts/cccl_preload/include/cub/util_namespace.cuh Normal file
View File

@@ -0,0 +1,172 @@
// SPDX-FileCopyrightText: Copyright (c) 2011, Duane Merrill. All rights reserved.
// SPDX-FileCopyrightText: Copyright (c) 2011-2021, NVIDIA CORPORATION. All rights reserved.
// SPDX-License-Identifier: BSD-3
/**
* \file util_namespace.cuh
* \brief Utilities that allow `cub::` to be placed inside an
* application-specific namespace.
*/
#pragma once
// This is not used by this file; this is a hack so that we can detect the
// CUB version from Thrust on older versions of CUB that did not have
// version.cuh.
#include <cub/version.cuh>
#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
# pragma GCC system_header
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
# pragma clang system_header
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
# pragma system_header
#endif // no system header
#include <cub/detail/detect_cuda_runtime.cuh>
// Prior to 1.13.1, only the PREFIX/POSTFIX macros were used. Notify users
// that they must now define the qualifier macro, too.
#if (defined(CUB_NS_PREFIX) || defined(CUB_NS_POSTFIX)) && !defined(CUB_NS_QUALIFIER)
# error CUB requires a definition of CUB_NS_QUALIFIER when CUB_NS_PREFIX/POSTFIX are defined.
#endif
#ifdef _CCCL_DOXYGEN_INVOKED
# define THRUST_CUB_WRAPPED_NAMESPACE
#endif
/**
* \def THRUST_CUB_WRAPPED_NAMESPACE
* If defined, this value will be used as the name of a namespace that wraps the
* `thrust::` and `cub::` namespaces.
* This macro should not be used with any other CUB namespace macros.
*/
#ifdef THRUST_CUB_WRAPPED_NAMESPACE
# define CUB_WRAPPED_NAMESPACE THRUST_CUB_WRAPPED_NAMESPACE
#endif
#ifdef _CCCL_DOXYGEN_INVOKED
# define CUB_WRAPPED_NAMESPACE
#endif
/**
* \def CUB_WRAPPED_NAMESPACE
* If defined, this value will be used as the name of a namespace that wraps the
* `cub::` namespace.
* If THRUST_CUB_WRAPPED_NAMESPACE is set, this will inherit that macro's value.
* This macro should not be used with any other CUB namespace macros.
*/
#ifdef CUB_WRAPPED_NAMESPACE
# define CUB_NS_PREFIX \
namespace CUB_WRAPPED_NAMESPACE \
{
# define CUB_NS_POSTFIX }
# define CUB_NS_QUALIFIER ::CUB_WRAPPED_NAMESPACE::cub
#endif
#ifdef _CCCL_DOXYGEN_INVOKED
# define CUB_NS_PREFIX
#endif
/**
* \def CUB_NS_PREFIX
* This macro is inserted prior to all `namespace cub { ... }` blocks. It is
* derived from CUB_WRAPPED_NAMESPACE, if set, and will be empty otherwise.
* It may be defined by users, in which case CUB_NS_PREFIX,
* CUB_NS_POSTFIX, and CUB_NS_QUALIFIER must all be set consistently.
*/
#ifndef CUB_NS_PREFIX
# define CUB_NS_PREFIX
#endif
#ifdef _CCCL_DOXYGEN_INVOKED
# define CUB_NS_POSTFIX
#endif
/**
* \def CUB_NS_POSTFIX
* This macro is inserted following the closing braces of all
* `namespace cub { ... }` block. It is defined appropriately when
* CUB_WRAPPED_NAMESPACE is set, and will be empty otherwise. It may be
* defined by users, in which case CUB_NS_PREFIX, CUB_NS_POSTFIX, and
* CUB_NS_QUALIFIER must all be set consistently.
*/
#ifndef CUB_NS_POSTFIX
# define CUB_NS_POSTFIX
#endif
#ifdef _CCCL_DOXYGEN_INVOKED
# define CUB_NS_QUALIFIER
#endif
/**
* \def CUB_NS_QUALIFIER
* This macro is used to qualify members of cub:: when accessing them from
* outside of their namespace. By default, this is just `::cub`, and will be
* set appropriately when CUB_WRAPPED_NAMESPACE is defined. This macro may be
* defined by users, in which case CUB_NS_PREFIX, CUB_NS_POSTFIX, and
* CUB_NS_QUALIFIER must all be set consistently.
*/
#ifndef CUB_NS_QUALIFIER
# define CUB_NS_QUALIFIER ::cub
#endif
#if defined(CUB_DISABLE_NAMESPACE_MAGIC) || defined(CUB_WRAPPED_NAMESPACE)
# if !defined(CUB_WRAPPED_NAMESPACE)
# if !defined(CUB_IGNORE_NAMESPACE_MAGIC_ERROR)
# error "Disabling namespace magic is unsafe without wrapping namespace"
# endif // !defined(CUB_IGNORE_NAMESPACE_MAGIC_ERROR)
# endif // !defined(CUB_WRAPPED_NAMESPACE)
# define CUB_DETAIL_MAGIC_NS_BEGIN
# define CUB_DETAIL_MAGIC_NS_END
#else // not defined(CUB_DISABLE_NAMESPACE_MAGIC)
# if defined(_NVHPC_CUDA)
# define CUB_DETAIL_MAGIC_NS_BEGIN \
inline namespace _CCCL_PP_CAT( \
_CCCL_PP_CAT(_CCCL_PP_CAT(_V_, CUB_VERSION), _CCCL_PP_SPLICE_WITH(_, _SM, NV_TARGET_SM_INTEGER_LIST)), _NVHPC) \
{
# define CUB_DETAIL_MAGIC_NS_END }
# else // not defined(_NVHPC_CUDA)
# define CUB_DETAIL_MAGIC_NS_BEGIN \
inline namespace _CCCL_PP_CAT(_CCCL_PP_CAT(_V_, CUB_VERSION), _CCCL_PP_SPLICE_WITH(_, _SM, __CUDA_ARCH_LIST__)) \
{
# define CUB_DETAIL_MAGIC_NS_END }
# endif // not defined(_NVHPC_CUDA)
#endif // not defined(CUB_DISABLE_NAMESPACE_MAGIC)
/**
* \def CUB_NAMESPACE_BEGIN
* This macro is used to open a `cub::` namespace block, along with any
* enclosing namespaces requested by CUB_WRAPPED_NAMESPACE, etc.
* This macro is defined by CUB and may not be overridden.
*/
#define CUB_NAMESPACE_BEGIN \
CUB_NS_PREFIX \
namespace cub \
{ \
CUB_DETAIL_MAGIC_NS_BEGIN
/**
* \def CUB_NAMESPACE_END
* This macro is used to close a `cub::` namespace block, along with any
* enclosing namespaces requested by CUB_WRAPPED_NAMESPACE, etc.
* This macro is defined by CUB and may not be overridden.
*/
#define CUB_NAMESPACE_END \
CUB_DETAIL_MAGIC_NS_END \
} /* end namespace cub */ \
CUB_NS_POSTFIX
// Declare these namespaces here for the purpose of Doxygenating them
CUB_NS_PREFIX
/*! \namespace cub
* \brief \p cub is the top-level namespace which contains all CUB
* functions and types.
*/
namespace cub
{
}
CUB_NS_POSTFIX

65
qwen3_6_scripts/cccl_preload/include/cub/version.cuh Normal file
View File

@@ -0,0 +1,65 @@
// SPDX-FileCopyrightText: Copyright (c) 2011-2022, NVIDIA CORPORATION. All rights reserved.
// SPDX-License-Identifier: BSD-3
/*! \file version.cuh
* \brief Compile-time macros encoding CUB release version
*
* <cub/version.h> is the only CUB header that is guaranteed to
* change with every CUB release.
*
*/
#pragma once
// For _CCCL_IMPLICIT_SYSTEM_HEADER
#include <cuda/__cccl_config>
#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
# pragma GCC system_header
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
# pragma clang system_header
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
# pragma system_header
#endif // no system header
#include <cuda/version>
/*! \def CUB_VERSION
* \brief The preprocessor macro \p CUB_VERSION encodes the version
* number of the CUB library as MMMmmmpp.
*
* \note CUB_VERSION is formatted as `MMMmmmpp`, which differs from `CCCL_VERSION` that uses `MMMmmmppp`.
*
* <tt>CUB_VERSION % 100</tt> is the sub-minor version.
* <tt>CUB_VERSION / 100 % 1000</tt> is the minor version.
* <tt>CUB_VERSION / 100000</tt> is the major version.
*/
#define CUB_VERSION 300500 // macro expansion with ## requires this to be a single value
/*! \def CUB_MAJOR_VERSION
* \brief The preprocessor macro \p CUB_MAJOR_VERSION encodes the
* major version number of the CUB library.
*/
#define CUB_MAJOR_VERSION (CUB_VERSION / 100000)
/*! \def CUB_MINOR_VERSION
* \brief The preprocessor macro \p CUB_MINOR_VERSION encodes the
* minor version number of the CUB library.
*/
#define CUB_MINOR_VERSION (CUB_VERSION / 100 % 1000)
/*! \def CUB_SUBMINOR_VERSION
* \brief The preprocessor macro \p CUB_SUBMINOR_VERSION encodes the
* sub-minor version number of the CUB library.
*/
#define CUB_SUBMINOR_VERSION (CUB_VERSION % 100)
/*! \def CUB_PATCH_NUMBER
* \brief The preprocessor macro \p CUB_PATCH_NUMBER encodes the
* patch number of the CUB library.
*/
#define CUB_PATCH_NUMBER 0
static_assert(CUB_MAJOR_VERSION == CCCL_MAJOR_VERSION);
static_assert(CUB_MINOR_VERSION == CCCL_MINOR_VERSION);
static_assert(CUB_SUBMINOR_VERSION == CCCL_PATCH_VERSION);