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 ....

[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.

[github-actions]

This branch currently is behind tinygrad/master. The line count difference bot is disabled.

[li0nr]

cleaned up and remvoed the matchers, tests are passing in both modes:
DEV=MOCK+CUDA:CUDA
DEV=MOCK+CUDA:PTX

[sirhcm]

can this work the same way MOCK+AMD::<arch> does?

[li0nr]

something like this ?

gpu_arch = int(MOCKGPU_ARCH[3:])
cuda_state.devices[ordinal] = {"compute_capability": divmod(gpu_arch, 10)}

i have added a commit

[geohot]

Are you confident the GPUOcelot port is complete, or did you just target making the tests pass?

[geohot]

Bounty locked to you, let's get this merged first. tinygrad/gpuocelot#12

[li0nr]

So my apprach was running tinygrad with the ocelot and list all the stuff we needed to support.
I will look at the MR there and address stuff.
Do we need/want to support everyting sm86 needs or just partial stuff?

[geohot]

We need everything in sm86! tinygrad may change and we need an emulator we can rely on