#include <cuda_runtime.h>
#include <cuda/std/bit>
#include <cuda/std/concepts>
#include <cstddef>
#include <cstdint>
#include <cusbf/cuda_error.hpp>
Go to the source code of this file.
|
| template<typename T > |
| __device__ __forceinline__ void | cusbf::detail::load256BitGlobalNC (const T *ptr, T *out) |
| | Loads 256 bits from global memory using the non-coherent cache path.
|
| |
| __device__ __forceinline__ void | cusbf::detail::load256BitGlobalNC (const uint64_t *ptr, uint64_t &out0, uint64_t &out1, uint64_t &out2, uint64_t &out3) |
| | Loads four uint64_t words via ld.global.nc.v4.u64 (sm_100+ only).
|
| |
| __device__ __forceinline__ void | cusbf::detail::load128BitGlobalNC (const uint64_t *ptr, uint64_t &out0, uint64_t &out1) |
| | Loads 128 bits from global memory using the non-coherent cache path.
|
| |
| __device__ __forceinline__ uint64_t | cusbf::detail::warpReduceOr (uint32_t peers, uint64_t value) |
| | OR-reduce a uint64_t across the lanes in a peer mask.
|
| |
| template<typename Kernel > |
| uint64_t | cusbf::detail::maxOccupancyGridSize (int32_t blockSize, Kernel kernel, uint64_t dynamicSMemSize) |
| | Calculates the maximum occupancy grid size for a kernel.
|
| |
