Cuda_nvrtc (sarek-cuda.Sarek_cuda.Cuda

Types

type nvrtc_program

NVRTC program handle

val nvrtc_program : nvrtc_program Ctypes.structure Ctypes.typ

val nvrtc_program_ptr : nvrtc_program Ctypes.structure Ctypes.ptr Ctypes.typ

type nvrtc_result =

| NVRTC_SUCCESS
| NVRTC_ERROR_OUT_OF_MEMORY
| NVRTC_ERROR_PROGRAM_CREATION_FAILURE
| NVRTC_ERROR_INVALID_INPUT
| NVRTC_ERROR_INVALID_PROGRAM
| NVRTC_ERROR_INVALID_OPTION
| NVRTC_ERROR_COMPILATION
| NVRTC_ERROR_BUILTIN_OPERATION_FAILURE
| NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION
| NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION
| NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID
| NVRTC_ERROR_INTERNAL_ERROR
| NVRTC_ERROR_UNKNOWN of int

NVRTC result codes

val nvrtc_result_of_int : int -> nvrtc_result

val int_of_nvrtc_result : nvrtc_result -> int

val nvrtc_result : nvrtc_result Ctypes.typ

val string_of_nvrtc_result : nvrtc_result -> string

Library Loading

val nvrtc_lib : Dl.library option Stdlib.Lazy.t

Load NVRTC library dynamically (lazy). Prefer unversioned to get system default that matches driver.

val is_available : unit -> bool

Check if NVRTC library is available

val get_nvrtc_lib : unit -> Dl.library

Get NVRTC library, raising if not available

val foreign_nvrtc_lazy : string -> ('a -> 'b) Ctypes.fn -> ('a -> 'b) lazy_t

Create a lazy foreign binding to NVRTC

Bindings

val nvrtcVersion_lazy : 
  (int Ctypes_static.ptr -> int Ctypes_static.ptr -> nvrtc_result) lazy_t

val nvrtcVersion : 
  int Ctypes_static.ptr ->
  int Ctypes_static.ptr ->
  nvrtc_result

val nvrtcCreateProgram_lazy : 
  (nvrtc_program Ctypes.structure Ctypes.ptr Ctypes_static.ptr ->
    string ->
    string option ->
    int ->
    string option Ctypes_static.ptr ->
    string option Ctypes_static.ptr ->
    nvrtc_result)
    lazy_t

val nvrtcCreateProgram : 
  nvrtc_program Ctypes.structure Ctypes.ptr Ctypes_static.ptr ->
  string ->
  string option ->
  int ->
  string option Ctypes_static.ptr ->
  string option Ctypes_static.ptr ->
  nvrtc_result

val nvrtcDestroyProgram_lazy : 
  (nvrtc_program Ctypes.structure Ctypes.ptr Ctypes_static.ptr ->
    nvrtc_result)
    lazy_t

val nvrtcDestroyProgram : 
  nvrtc_program Ctypes.structure Ctypes.ptr Ctypes_static.ptr ->
  nvrtc_result

val nvrtcCompileProgram_lazy : 
  (nvrtc_program Ctypes.structure Ctypes.ptr ->
    int ->
    string Ctypes_static.ptr ->
    nvrtc_result)
    lazy_t

val nvrtcCompileProgram : 
  nvrtc_program Ctypes.structure Ctypes.ptr ->
  int ->
  string Ctypes_static.ptr ->
  nvrtc_result

val nvrtcGetPTXSize_lazy : 
  (nvrtc_program Ctypes.structure Ctypes.ptr ->
    Unsigned.size_t Ctypes_static.ptr ->
    nvrtc_result)
    lazy_t

val nvrtcGetPTXSize : 
  nvrtc_program Ctypes.structure Ctypes.ptr ->
  Unsigned.size_t Ctypes_static.ptr ->
  nvrtc_result

val nvrtcGetPTX_lazy : 
  (nvrtc_program Ctypes.structure Ctypes.ptr ->
    char Ctypes_static.ptr ->
    nvrtc_result)
    lazy_t

val nvrtcGetPTX : 
  nvrtc_program Ctypes.structure Ctypes.ptr ->
  char Ctypes_static.ptr ->
  nvrtc_result

val nvrtcGetCUBINSize_lazy : 
  (nvrtc_program Ctypes.structure Ctypes.ptr ->
    Unsigned.size_t Ctypes_static.ptr ->
    nvrtc_result)
    option
    lazy_t

val nvrtcGetCUBINSize : 
  nvrtc_program Ctypes.structure Ctypes.ptr ->
  Unsigned.size_t Ctypes_static.ptr ->
  nvrtc_result

val nvrtcGetCUBIN_lazy : 
  (nvrtc_program Ctypes.structure Ctypes.ptr ->
    char Ctypes_static.ptr ->
    nvrtc_result)
    option
    lazy_t

val nvrtcGetCUBIN : 
  nvrtc_program Ctypes.structure Ctypes.ptr ->
  char Ctypes_static.ptr ->
  nvrtc_result

val nvrtcGetProgramLogSize_lazy : 
  (nvrtc_program Ctypes.structure Ctypes.ptr ->
    Unsigned.size_t Ctypes_static.ptr ->
    nvrtc_result)
    lazy_t

val nvrtcGetProgramLogSize : 
  nvrtc_program Ctypes.structure Ctypes.ptr ->
  Unsigned.size_t Ctypes_static.ptr ->
  nvrtc_result

val nvrtcGetProgramLog_lazy : 
  (nvrtc_program Ctypes.structure Ctypes.ptr ->
    char Ctypes_static.ptr ->
    nvrtc_result)
    lazy_t

val nvrtcGetProgramLog : 
  nvrtc_program Ctypes.structure Ctypes.ptr ->
  char Ctypes_static.ptr ->
  nvrtc_result

val nvrtcAddNameExpression_lazy : 
  (nvrtc_program Ctypes.structure Ctypes.ptr -> string -> nvrtc_result) lazy_t

val nvrtcAddNameExpression : 
  nvrtc_program Ctypes.structure Ctypes.ptr ->
  string ->
  nvrtc_result

val nvrtcGetLoweredName_lazy : 
  (nvrtc_program Ctypes.structure Ctypes.ptr ->
    string ->
    string Ctypes_static.ptr ->
    nvrtc_result)
    lazy_t

val nvrtcGetLoweredName : 
  nvrtc_program Ctypes.structure Ctypes.ptr ->
  string ->
  string Ctypes_static.ptr ->
  nvrtc_result

High-Level Helpers

exception Nvrtc_error of nvrtc_result * string

Exception for NVRTC errors

val check : string -> nvrtc_result -> unit

Check result and raise if error

val cuda_target_triples : string list

CUDA header search path

NVRTC has NO default include path. It is a library, not a driver, so it does not inherit the built-in -I that nvcc adds, and __half is not an nvrtc builtin. A generated kernel that says #include <cuda_fp16.h> — which is exactly what the f16 codegen emits, see Sarek_codegen.Sarek_ir_cuda.cuda_fp16_include — therefore fails with NVRTC_ERROR_COMPILATION and

could not open source file "cuda_fp16.h" (no directories in search list)

unless the toolkit's include directory is passed explicitly. Verified against libnvrtc 13.3: byte-identical source compiles to PTX containing cvt.rn.f16.f32 with the flag and fails without it.

Discovery order (every surviving candidate is passed, so a partial toolkit layout still resolves):

SAREK_CUDA_INCLUDE — explicit ':'-separated override, kept if the directory exists;
CUDA_PATH / CUDA_HOME / CUDA_ROOT derived include and targets/<triple>/include;
a short list of conventional install roots.

Derived and conventional candidates must additionally CONTAIN cuda_fp16.h. That marker is what keeps the discovery honest: a stale CUDA_PATH or a bare /usr/include never becomes an -I that could shadow a real header. On a machine with no CUDA headers the list is empty and the option array is byte-identical to before, so non-CUDA hosts are unaffected.

Passed unconditionally rather than gated on the f16 detector: this module never sees the IR, an unused include directory cannot change a kernel that includes nothing, and gating here would leave any future header-using codegen broken in the same way.

val cuda_conventional_roots : string list

val cuda_fp16_header : string

val is_dir : string -> bool

val has_fp16_header : string -> bool

val env_nonempty : string -> string option

val cuda_include_paths : string list Stdlib.Lazy.t

Existing CUDA include directories, most-specific first, de-duplicated.

val cuda_include_flags : string list Stdlib.Lazy.t

The --include-path= flags derived from cuda_include_paths.

exception Fp_conformance_violation of string

Floating-point conformance guard

Raised when an nvrtc option would change binary32 semantics in a way the interpreter — Sarek's cross-backend oracle — cannot follow. Carries the offending option and why it is refused.

The floating-point option classes this path refuses, and why.

The rule these encode: Sarek's DSL semantics are IEEE-754 binary32 with every operation rounded as written, and the interpreter implements exactly that. An option that FLUSHES SUBNORMALS or DOWNGRADES a division/square root to an approximate form makes the device disagree with the interpreter on inputs the test suite already uses, and — unlike HIP's contraction case, where appending -ffp-contract=off last neutralises the caller — there is no later flag that undoes it. So these are rejected rather than warned about.

-use_fast_math implies --ftz=true --prec-div=false --prec-sqrt=false --fmad=true (CUDA nvrtc documentation, dependencies/Cuda/nvrtc.h).
-ftz=true flushes binary32 subnormals to zero. MEASURED, host-side on CUDA 13.3 (nvcc/ptxas V13.3.73, no NVIDIA device): compiling the generated f16_midround kernel for sm_90 with -use_fast_math or -ftz=true turns FMUL/FADD into FMUL.FTZ/FADD.FTZ. 1e-5 is already lane 5 of the test_hip_f16 input array, so this is reachable from the existing test data, not a hypothetical.
--prec-div=false / --prec-sqrt=false select approximate division and square root. Sarek_df64 already spends its whole error budget on the correctly-rounded forms — the PTX emitter switched div.approx.f32 to div.rn.f32 for exactly this reason (audit finding M2).

--fmad=true is NOT refused: it is nvrtc's default, so refusing it would refuse the status quo. Contraction is instead defeated by construction where it matters (see Sarek_df64's contraction barrier). It warns, so a caller who sets it deliberately is told the guarantee it is trading away.

NVRTC ACCEPTS AN OPTION AND ITS VALUE AS TWO SEPARATE ARRAY ELEMENTS. An earlier version of this guard matched ["--ftz=true"] and walked straight past ["--ftz"; "true"] — executed against libnvrtc 13.3 through these bindings, the separated form compiled and the PTX carried .ftz on the f32 arithmetic, with no exception and no warning. So the scan below is written around the OPTION, not its spelling: a value-taking name consumes the next array element when there is no =, and is FAIL-CLOSED — --ftz with no determinable value is refused, because a guard that cannot tell what a flag is set to must not assume the safe answer. (Hip_rtc.fp_relaxing_option_prefixes never had this hole; it matches by prefix on options whose value is always inline.)

type fp_verdict =

| Fp_reject of string
| Fp_warn of string

val fp_option_classes : 
  (string * (string option -> bool) * [ `Reject | `Warn ] * string) list

(name, value_semantics, why). value_semantics answers "is this setting dangerous?" given the resolved value, where None means the value could not be determined.

val fp_value_taking_names : string list

Names that take a value, and may therefore carry it in the NEXT array element. use_fast_math is a bare switch and is deliberately absent.

val split_nvrtc_option : string -> string * string option

Split an nvrtc option into its name (leading dashes stripped) and inline value, if any. --ftz=true -> ("ftz", Some "true"); -use_fast_math -> ("use_fast_math", None).

val looks_like_option : string -> bool

val fp_scan : string list -> fp_verdict list

Scan a WHOLE option array, resolving values that live in a following element. Pure, so it can be exercised without libnvrtc or a device — and list-level, because the separated form is invisible to any per-element check.

val fp_rejection_reason_list : string list -> string option

The reason the option array opts must be refused, or None.

val fp_rejection_reason : string -> string option

Single-element convenience wrapper. Correct only for options that carry their value inline — use fp_rejection_reason_list for anything that might be split across elements.

val fp_warning_reason_list : string list -> string option

val fp_warning_reason : string -> string option

val check_fp_conformance : string list -> unit

Reject FP-relaxing options and warn about the merely risky ones.

Applied to the WHOLE array, at the single point where an array reaches nvrtcCompileProgram, so it covers options from any source — a caller, a future environment-variable escape hatch, or a hardcoded flag added by a later maintainer — and both spellings of every value-taking option.

val nvrtc_option_array : 
  ?arch:string ->
  options:string list ->
  include_opts:string list ->
  unit ->
  string list

The exact option array this module hands to nvrtcCompileProgram for one architecture attempt. Split out so the composition (caller options + the module's own flags) can be screened in a test, not only the caller's half — a hardcoded FP flag added here later must be caught the same way a caller's is.

val compile_to_ptx : 
  ?name:string ->
  arch:string ->
  ?options:string list ->
  string ->
  string

Compile CUDA source to PTX.

parameter source
CUDA C source code

parameter name
Optional program name

parameter arch
Target architecture (e.g., "compute_75")

parameter options
Extra nvrtc options, prepended to the ones this module supplies. Screened by check_fp_conformance: an option that would break binary32 agreement with the interpreter raises Fp_conformance_violation BEFORE any nvrtc call, so the check is reachable on a host with no CUDA at all.

returns
PTX code as string

val get_version : unit -> int * int

Get NVRTC version as (major, minor)