Cuda 12 sm86

[li0nr]

supporting cuda 12 and sm_86 in ocelot.
bf16/f16 ops and mma ops added

[geohot]

So it's not complete. For example, min.f16 missing while max.f16 is implemented.

Also, here's a GPT-6 review, I think it has good points.

Review: request changes

Reviewed commit f0c3690. Four correctness issues should be addressed before merging.

1. [P2] TF32 conversion is missing, preventing normal TF32 MMA kernels from loading

Location: ocelot/src/parser/ptxgrammar.yy:307–310 (https://github.com/tinygrad/gpuocelot/blob/f0c3690f/ocelot/src/parser/ptxgrammar.yy#L307-L310)

The PR accepts TF32 MMA instructions, but not the SM86 conversion used to prepare their inputs:

  .reg .f32 a;
  .reg .b32 t;
  cvt.rna.tf32.f32 t, a;

NVIDIA ptxas accepts this; Ocelot rejects it. TF32 is absent from conversion datatypes, .rna is unrecognized, and eval_Cvt lacks the conversion.

Rounding inside eval_Mma cannot compensate: kernels containing the conversion fail before execution. Implement the conversion end-to-end, including its ties-away rounding semantics.

For SM86, the relevant instruction is .rna, not .rn; NVIDIA requires SM90 and PTX 7.8 for cvt.rn.tf32.f32.

2. [P2] FP16 arithmetic silently ignores .sat

Location: ocelot/src/executive/CooperativeThreadArray.cpp:1783–1784 (https://github.com/tinygrad/gpuocelot/blob/f0c3690f/ocelot/src/executive/CooperativeThreadArray.cpp#L1783-L1784)

The new FP16 arithmetic branches pass the modifier to toF16, but that helper only handles rounding—not saturation.

Confirmed:

  add.rn.sat.f16(1, 1)
  Expected: 1  (0x3c00)
  Actual:   2  (0x4000)

The new subtraction, multiplication, and FP16 FMA paths likewise omit saturation. Apply PTX’s [0,1] saturation semantics, including NaN-to-zero behavior, and add modifier-specific
tests.

3. [P2] FP16 FMA produces incorrect results through double rounding

Location: ocelot/src/executive/CooperativeThreadArray.cpp:4648 (https://github.com/tinygrad/gpuocelot/blob/f0c3690f/ocelot/src/executive/CooperativeThreadArray.cpp#L4648)

Because a, b, and c are FP32, std::fma(a,b,c) rounds to FP32 before toF16 rounds again. This is not equivalent to one correctly rounded FP16 fused operation.

Confirmed for FP16 input bit patterns:

  a = 0x3e00, b = 0x3956, c = 0x0001

  Expected fma.rn.f16: 0x3c01
  Actual:             0x3c00

The small positive addend is lost during intermediate FP32 rounding. Avoid intermediate rounding that can change the final FP16 result, and include this case as a regression test.

4. [P2] FP16 FMA flushes subnormal inputs but not outputs

Location: ocelot/src/executive/CooperativeThreadArray.cpp:4647–4648 (https://github.com/tinygrad/gpuocelot/blob/f0c3690f/ocelot/src/executive/CooperativeThreadArray.cpp#L4647-L4648)

The FP16 FMA branch applies ftzF16 to each input but stores the converted result without flushing it. PTX .ftz applies to subnormal results too.

Confirmed:

  fma.rn.ftz.f16(2^-14, 0.5, 0)
  Expected: 0x0000
  Actual:   0x0200

Apply ftzF16 to the converted output, preserving the sign of zero, as the new add/sub/mul branches already do.