chunk_geometry.h

Go to the documentation of this file.

#pragma once
 
namespace fz {
namespace fused {
namespace chunk {
 
constexpr int TEMP_BYTES = 4096;   // LC coder scratch (RZE/RRE/RARE/RAZE/CLOG/HCLOG all 4096)
constexpr int TPB        = 512;    // = LC coder TPB
 
constexpr int kSupportedChunkBytes[] = {4096, 8192, 16384};
constexpr int kNumSupportedChunkBytes =
    sizeof(kSupportedChunkBytes) / sizeof(kSupportedChunkBytes[0]);
 
// Host-only (launcher validation). Excluded from NVRTC's device-JIT compile
// (__CUDACC_RTC__, defined automatically by NVRTC) via #ifndef below: NVRTC
// rejects any unannotated function merely being *present* in the translation
// unit in JIT mode (not just called), and this header can't use
// __host__/__device__ either, since it's also included from plain host C++
// translation units where those tokens aren't defined at all (no CUDA
// compiler). Geom<>'s static_assert below inlines the same check as a literal
// expression instead of calling this, so it works unmodified under NVRTC too.
#ifndef __CUDACC_RTC__
constexpr bool isSupportedChunkBytes(int bytes) {
    return bytes == 4096 || bytes == 8192 || bytes == 16384;
}
#endif
 
template <int Bytes>
struct Geom {
    static_assert(Bytes == 4096 || Bytes == 8192 || Bytes == 16384,
                 "chunk-cooperative chunk size must be one of kSupportedChunkBytes");
    static constexpr int CHUNK_BYTES = Bytes;
    static constexpr int NELEM       = Bytes / 4;   // uint32 codes per chunk
    static constexpr int NPP         = NELEM / 32;  // bitshuffle words per plane
};
 
} // namespace chunk
} // namespace fused
} // namespace fz