helpers.cuh

Go to the documentation of this file.

#pragma once
 
#include <cuda_runtime.h>
 
#include <cuda/std/bit>
#include <cuda/std/concepts>
 
#include <cstddef>
#include <cstdint>
 
#include <cusbf/cuda_error.hpp>
 
namespace cusbf::detail {
 
#if __CUDA_ARCH__ >= 1000
 
template <typename T>
__device__ __forceinline__ void load256BitGlobalNC(const T* ptr, T* out) {
    static_assert(sizeof(T) == 4 || sizeof(T) == 8, "T must be uint32_t or uint64_t");
 
    if constexpr (sizeof(T) == 8) {
        asm volatile("ld.global.nc.v4.u64 {%0, %1, %2, %3}, [%4];"
                     : "=l"(out[0]), "=l"(out[1]), "=l"(out[2]), "=l"(out[3])
                     : "l"(ptr));
    } else {
        asm volatile("ld.global.nc.v8.u32 {%0, %1, %2, %3, %4, %5, %6, %7}, [%8];"
                     : "=r"(out[0]),
                       "=r"(out[1]),
                       "=r"(out[2]),
                       "=r"(out[3]),
                       "=r"(out[4]),
                       "=r"(out[5]),
                       "=r"(out[6]),
                       "=r"(out[7])
                     : "l"(ptr));
    }
}
 
__device__ __forceinline__ void load256BitGlobalNC(
    const uint64_t* ptr,
    uint64_t& out0,
    uint64_t& out1,
    uint64_t& out2,
    uint64_t& out3
) {
    asm volatile("ld.global.nc.v4.u64 {%0, %1, %2, %3}, [%4];"
                 : "=l"(out0), "=l"(out1), "=l"(out2), "=l"(out3)
                 : "l"(ptr));
}
 
#endif
 
__device__ __forceinline__ void
load128BitGlobalNC(const uint64_t* ptr, uint64_t& out0, uint64_t& out1) {
    asm volatile("ld.global.nc.v2.u64 {%0, %1}, [%2];" : "=l"(out0), "=l"(out1) : "l"(ptr));
}
 
__device__ __forceinline__ uint64_t warpReduceOr(uint32_t peers, uint64_t value) {
#if __CUDA_ARCH__ >= 800
    auto lo = __reduce_or_sync(peers, static_cast<uint32_t>(value));
    auto hi = __reduce_or_sync(peers, static_cast<uint32_t>(value >> 32));
    return (static_cast<uint64_t>(hi) << 32) | lo;
#else
    // Shuffle-based reduction across the lanes set in `peers`.
    uint32_t remaining = peers;
    while (remaining) {
        int src = __ffs(remaining) - 1;
        uint64_t other =
            (static_cast<uint64_t>(__shfl_sync(peers, static_cast<uint32_t>(value >> 32), src))
             << 32) |
            __shfl_sync(peers, static_cast<uint32_t>(value), src);
        value |= other;
        remaining &= remaining - 1;  // clear lowest set bit
    }
    return value;
#endif
}
 
template <typename Kernel>
uint64_t maxOccupancyGridSize(int32_t blockSize, Kernel kernel, uint64_t dynamicSMemSize) {
    int device = 0;
    cudaGetDevice(&device);
 
    int numSM = -1;
    cudaDeviceGetAttribute(&numSM, cudaDevAttrMultiProcessorCount, device);
 
    int maxActiveBlocksPerSM{};
    cudaOccupancyMaxActiveBlocksPerMultiprocessor(
        &maxActiveBlocksPerSM, kernel, blockSize, dynamicSMemSize
    );
 
    return maxActiveBlocksPerSM * numSM;
}
 
}  // namespace cusbf::detail