gpuocelot cuda12 sm 86

[li0nr]

This PR uses the GPUOcelot CUDA 12 sm86 release (which is the main change), adds CUDA 12 NVRTC/runtime header setup.

Gpuocelot add support to cuda12 and sm_50 and sm_86.

at the beginng i just wanted for gpuocelot to run with cuda-12 with minimal arch support sm_50.
Changes where small we can see it in this branch https://github.com/li0nr/tinygrad/commits/gpuocelot-cuda-12.
if it is safer we can go this route.

The next best thing (to be bounty worthy) was supporting sm_86 arch in gpuocelot.

this needs to be modfied after we merge the ocelot branch to the tinygrad repo.
tinygrad/gpuocelot#12 / https://github.com/li0nr/gpuocelot/tree/cuda-12-sm86

CI: li0nr#2

[li0nr]

I neede to add this bc of failures like this https://github.com/li0nr/tinygrad/actions/runs/33324336717/job/99291747648,
maybe bc of this condition if (nvrtcMajor.value, nvrtcMinor.value) >= (12, 4): self.compile_options.append("--minimal")
If there is any cleaner way of doing this, I can change

[li0nr]

If we approve the changes to ocelot, I will change this to be more offical release ....

[github-actions]

Changes

Name                           Lines    Diff    Tokens/Line    Diff
---------------------------  -------  ------  -------------  ------
tinygrad/renderer/cstyle.py      502     +10           20.5    +0.2


total lines changes: +10

[geohot]

Tests passing? You need to remove those matchers from cstyle, it needs to work on 100% of SM86 PTX.

[li0nr]

yep tests are passing.

the matchers are for sm86, the cuda cpp generated code is being rejected by the new nvrtc, for example :

  *(data0_16 + alu0) = ((long)(alu3));

  Exact NVRTC error:

  <null>(26): error: more than one conversion function from "half" (aka "__half") to "long" applies:
      function "__half::operator float() const"
      function "__half::operator short() const"
      function "__half::operator unsigned short() const"
      function "__half::operator int() const"
      function "__half::operator unsigned int() const"
      function "__half::operator long long() const"
      function "__half::operator unsigned long long() const"
      function "__half::operator __nv_bool() const"

  *(data0_16+(alu0+1)) = ((long)(alu6));
                                ^
  4 errors detected in the compilation of "<null>".

or maybe i did not understand your point.

running the full test suite with sm86 with these matchers is passing.

[geohot]

You don't need the matchers for 3090/4090/5090, what version is that?

[li0nr]

I chose the SM86 specifaclly because it matches 3090 according to this https://developer.nvidia.com/cuda/gpus.
I will try to targeting 4090/5090 sm. As well as testing cuda11 with sm_86.

[li0nr]

found the issue, I was not using the correct cuda_cudart, in my linux env.
I will clean the matcher.