//===-- cstdint - the six integer widths FlashInfer counts in ------------===//
//
// `uint32_t` and its five relatives, and the two `*_C` literal macros
// `fp16.h` spells its constants with. This is what `#include <cstdint>`
// resolves to when the compiler is NVRTC and the include path is a header set
// carried in the binary rather than a directory on a disk.
//
// # Why this is not a guard
//
// Every other host header in FlashInfer's attention closure -- `<iostream>`,
// `<sstream>`, `<vector>`, `<cuda_runtime.h>` and eighteen more -- is reached
// only by host code, so the vendored copy guards the directive with
// `#ifndef __CUDACC_RTC__` and the names go away with the code that wanted
// them. `<cstdint>` is different, and the difference was measured rather than
// assumed: `uint32_t` appears **2512 times** in the closure, `uint8_t` 194,
// `int64_t` 127, `uint16_t` 47, `int32_t` 37, `uint64_t` 11 -- overwhelmingly
// in `__device__` templates and `__global__` kernels. Guarding the include
// away leaves every one of them undefined.
//
// And NVRTC supplies none of them. Measured on this box (L40S, NVRTC 13.0),
// with the real `cuda_fp16.h` / `cuda_bf16.h` / `cuda_fp8.h` / `cuda_fp4.h`
// in the header set:
//
// ```text
// __global__ void probe(unsigned* p) { uint32_t x = 1; *p = x; }
//   -> error: identifier "uint32_t" is undefined
// ```
//
// NVRTC's built-in preamble has `size_t`, `ptrdiff_t`, `nullptr_t`, `assert`,
// `printf`, `memcpy` and the `f`-suffixed math functions. It has no
// fixed-width types, no `INT32_MAX`, no `FLT_MAX`, and no `std::` anything.
// So this file exists, and it is the whole of `<cstdint>` this closure uses.
//
// # The name is the trick, as it is for the CCCL shims
//
// NVRTC resolves an `#include` by matching the directive's LITERAL spelling
// against `includeNames[]`, so an entry named `cstdint` answers
// `#include <cstdint>` from source we do not own, leaving that source
// unmodified. `csrc/shim/cooperative_groups.h` makes the same argument at
// greater length; this is the same rule applied to the C++ standard library
// instead of to CCCL.
//
// # What is deliberately NOT here
//
// The macros. `INT32_MAX`, `SIZE_MAX`, `INTMAX_C` and the rest of
// `<stdint.h>`'s constant surface are absent because the closure does not name
// one of them -- a grep over all 28 files found `UINT32_C` and `UINT16_C` and
// nothing else. A shim that grew to the standard's surface would be a second
// set of promises about integer formats to disagree with the first, and the
// disagreement would be silent.
//
// `std::uint32_t`. A real `<cstdint>` puts the names in `std` and then drags
// them into the global namespace; FlashInfer spells them unqualified at all
// 2928 sites, so the `std` half would be dead text.
//
// # The `std` half, added later, because that premise stopped being true
//
// It was dead text while FlashInfer was the only includer. It is not any more:
// `csrc/src` -- OUR headers, which nvcc compiles today and NVRTC compiles
// under the JIT -- spells `std::uint32_t` **286** times, `std::uint64_t` 102,
// `std::uint8_t` 41, `std::int32_t` 16, `std::int64_t` 6 and `std::uintptr_t`
// 3. The FA2 lattice's root reaches two of those files and the compile stops
// with *"namespace `std` has no member `uint32_t`"* at ten sites.
//
// So the `std` half is here, and the paragraph above is kept rather than
// deleted because its ARGUMENT is still the right one -- carry the names the
// actual includers write, and no more. What changed is the includer set, not
// the rule, and the counts are what changed it. Found with a `/tmp` `libnvrtc`
// probe against `csrc/src/attn/fa2.cuh`, before any of it had been built.
//
// The names are re-exported, not re-declared. `using ::uint32_t;` makes
// `std::uint32_t` and `::uint32_t` the same type by construction; a second
// `typedef` could be edited into a second opinion, and a disagreement between
// the two spellings surfaces as a template deduction failure a hundred frames
// deep rather than as an error about integer widths. The two literal macros
// stay global-only -- a macro has no namespace, and `std::UINT32_C` is not a
// thing anyone can write.

#ifndef PIE_NVRTC_CSTDINT_
#define PIE_NVRTC_CSTDINT_

// NVRTC's preamble already declares `size_t` and `ptrdiff_t`; naming them
// again here would be a redefinition, not a courtesy.
typedef signed char int8_t;
typedef short int16_t;
typedef int int32_t;
typedef long long int64_t;

typedef unsigned char uint8_t;
typedef unsigned short uint16_t;
typedef unsigned int uint32_t;
typedef unsigned long long uint64_t;

// The two pointer-width integers. `uintptr_t` reaches device code through
// `FLASHINFER_CHECK_ALIGNMENT` and through the `reinterpret_cast<uintptr_t>`
// that `cp_async.cuh` and `permuted_smem.cuh` use to turn a shared-memory
// address into the 32-bit form `cp.async` wants; 64-bit is the only device
// ABI this crate targets.
//
// `long long` and not `long`, which is not a style choice: CCCL's own
// `cuda/std/cstdint` declares these two in the GLOBAL namespace under
// `__CUDACC_RTC__`, as `int64_t`/`uint64_t`, and a header set that carries
// both files gets *"invalid redeclaration of type name `intptr_t`"* if the
// two spellings pick different underlying types. Same width, same ABI, one
// declaration.
typedef int64_t intptr_t;
typedef uint64_t uintptr_t;

// `fp16.h` writes its IEEE constants as `UINT32_C(0x77800000)`. The suffix
// has to make the literal the width the typedef above is, which on this ABI
// is `unsigned int` and `unsigned short` promoted -- so `U` and nothing more.
#define UINT16_C(value) value##U
#define UINT32_C(value) value##U

// The `std` half. See the header comment: `csrc/src` writes these qualified at
// 454 sites and the vendored closure writes them unqualified at 2928, so both
// spellings have to name the same six types.
namespace std {

using ::int16_t;
using ::int32_t;
using ::int64_t;
using ::int8_t;

using ::uint16_t;
using ::uint32_t;
using ::uint64_t;
using ::uint8_t;

using ::intptr_t;
using ::uintptr_t;

}  // namespace std

#endif  // PIE_NVRTC_CSTDINT_
