ckitching · HackerTrans

2 वर्ष पहलेdiscuss

The LLVM manual touches on some of the basics of why: https://llvm.org/docs/CompileCudaWithLLVM.html#dialect-diffe...

2 वर्ष पहलेdiscuss

Not any more ;)

2 वर्ष पहलेdiscuss

Prettymuch. Compilers can do a lot more than people give them credit for. At least AMD document their hardware so it is actually possible to know low-level details. PTX can obfuscate that surprisingly badly for nvidia targets.

ckitching

2 वर्ष पहलेdiscuss

Oh yes, we found all kinds of bugs in Nvidia's cuda implementation during this project :D.

There's a bunch of pretty obscure functions in the device side apis too: some esoteric math functions, old simd "intrinsics" that are mostly irrelevant with modern compilers, etc.

ckitching

2 वर्ष पहलेdiscuss

SCALE doesn't use cuBlas and friends. For those APIs, it uses either its own implementations of the functions, or delegates to an existing AMD library (such as rocblas).

It wouldn't even be technically possible for SCALE to distribute and use cuBlas, since the source code is not available. I suppose maybe you could do distribute cuBlas and run it through ZLUDA, but that would likely become legally troublesome.

ckitching

2 वर्ष पहलेdiscuss

The CUDA C APIs are supported as much in C as in C++ using SCALE!

Cuda-fortran is not currently supported by scale since we haven't seen much use of it "in the wild" to push it up our priority list.

ckitching

2 वर्ष पहलेdiscuss

Indeed, no extra synchronisation is needed here due to the nature of the hardware (threads in a warp can't get out of sync with each other).

Even on NVIDIA, you could've written this without the asm a discussed above!

ckitching

2 वर्ष पहलेdiscuss

I certainly would not attempt this feat with x86 `asm` blocks :D. PTX is indeed very pedestrian: it's more like IR than machine code, really. All the usual "machine-level craziness" that would otherwise make this impossible is just unrepresentable in PTX (though you do run into cases of "oopsie, AMD don't have hardware for this so we have to do something insane").

ckitching

2 वर्ष पहलेdiscuss

NVCC is doing much better than before in terms of "broken C++". There was indeed a time when lots of modern C++ just didn't work.

Nowadays the issues are more subtle and nasty. Subtle differences in overload resolution. Subtle differences in lambda handling. Enough to break code in "spicy" ways when you try to port it over.

ckitching

2 वर्ष पहलेdiscuss

Add one: it's trivial to add a compiler builtin to carry the instruction from the frontend to the backend if an instruction exists and the backend knows about it.

If there's no instruction, either, you can write a C++ function to replicate the behaviour and codegen a call to it. Since the PTX blocks are expanded during initial IR generation, it all inlines nicely by the end. Of course, such software emulation is potentially suboptimal (depends on the situation).

ckitching

2 वर्ष पहलेdiscuss

I compiled your function with SCALE for gfx1030:

        .p2align        2                               ; -- Begin function _Z15ptx_thread_voteff
        .type   _Z15ptx_thread_voteff,@function
  _Z15ptx_thread_voteff:                  ; @_Z15ptx_thread_voteff
  ; %bb.0:                                ; %entry
        s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
        s_waitcnt_vscnt null, 0x0
        v_cmp_ge_f32_e32 vcc_lo, v0, v1
        s_cmp_eq_u32 vcc_lo, -1
        s_cselect_b32 s4, -1, 0
        v_cndmask_b32_e64 v0, 0, 1, s4
        s_setpc_b64 s[30:31]
  .Lfunc_end1:
        .size   _Z15ptx_thread_voteff, .Lfunc_end1-_Z15ptx_thread_voteff
                                        ; -- End function

What were the safety concerns you had? This code seems to be something like `return __all_sync(rSq >= rCritSq) ? 1 : 0`, right?

ckitching

2 वर्ष पहलेdiscuss

Rather awkwardly, you've asked about an instruction that isn't currently implemented. :D Support for wmma and friends is in development.

But in general the answer to your question is yes: we use AMD-specific builtins where available/efficient to make things work. Otherwise many things would be unrepresentble, not just slow!

ckitching

2 वर्ष पहलेdiscuss

> Hi, why do you believe that bfloat16 is not supported?

Apologies, I appear to be talking nonsense. I conflated bfloat16 with nvidia's other wacky floating point formats. This is probably my cue to stop answering reddit/HN comments and go to bed. :D

So: ahem: bfloat16 support is basically just missing the fairly boring header.

> Regarding cublasLt, what is your plan for support there? Pass everything through to hipblasLt (hipify style) or something else?

Prettymuch that, yes. Not much point reimplementing all the math libraries when AMD is doing that part of the legwork already.

ckitching

2 वर्ष पहलेdiscuss

> You might also find raw c++ for device libraries saner to deal with than cuda.

Actually, we just compile all the device libraries to LLVM bitcode and be done with it. Then we can write them using all the clang-dialect, not-nvcc-emulating, C++23 we feel like, and it'll still work when someone imports them into their c++98 CUDA project from hell. :D

ckitching

2 वर्ष पहलेdiscuss

Greetings, I work on SCALE.

It appears we implemented `--threads` but not `-t` for the compiler flag. Oeps. In either case, the flag has no effect at present, since fatbinary support is still in development, and that's the only part of the process that could conceivably be parallelised.

That said: clang (and hence the SCALE compiler) tends to compile CUDA much faster than nvcc does, so this lack of the parallelism feature is less problematic than it might at first seem.

NVTX support (if you want more than just "no-ops to make the code compile") requires cooperation with the authors of profilers etc., which has not so far been available

bfloat16 is not properly supported by AMD anyway: the hardware doesn't do it, and HIP's implementatoin just lies and does the math in `float`. For that reason we haven't prioritised putting together the API.

cublasLt is a fair cop. We've got a ticket :D.

ckitching

2 वर्ष पहलेdiscuss

It works exactly as well as other AMDGPU-related software (HIP etc.) works inside Docker.

There are some delightful AMD driver issues that make certain models of GPU intermittently freeze the kernel when used from docker. That was great fun when building SCALE's CI system :D.

ckitching

2 वर्ष पहलेdiscuss

[I work on SCALE]

CUDA has a couple of extra problems beyond just any other programming language:

- CUDA is more than a language: it's a giant library (for both CPU and GPU) for interacting with the GPU, and for writing the GPU code. This needed reimplementing. At least for the device-side stuff we can implement it in CUDA, so when we add support for other GPU vendors the code can (mostly) just be recompiled and work there :D. - CUDA (the language) is not actually specified. It is, informally, "whatever nvcc does". This differs significantly from what Clang's CUDA support does (which is ultimately what the HIP compiler is derived from).

PTX is indeed vastly annoying.

ckitching

2 वर्ष पहलेdiscuss

Like this:

https://docs.scale-lang.com/manual/how-to-use/#identifying-g...

ckitching

2 वर्ष पहलेdiscuss

You're right that most people only use a small subset of cuda: we prioritied support for features based on what was needed for various open-source projects, as a way to try to capture the most common things first.

A complete API comparison table is coming soon, I belive. :D

In a nutshell: - DPX: Yes. - Shuffles: Yes. Including the PTX versions, with all their weird/wacky/insane arguments. - Atomics: yes, except the 128-bit atomics nvidia added very recently. - MMA: in development, though of course we can't fix the fact that nvidia's hardware in this area is just better than AMD's, so don't expect performance to be as good in all cases. - TMA: On the same branch as MMA, though it'll just be using AMD's async copy instructions.

> mapping every PTX instruction to a direct RDNA counterpart or a list of instructions used to emulate it.

We plan to publish a compatibility table of which instructons are supported, but a list of the instructions used to produce each PTX instruction is not in general meaningful. The inline PTX handler works by converting the PTX block to LLVM IR at the start of compilation (at the same time the rest of your code gets turned into IR), so it then "compiles forward" with the rest of the program. As a result, the actual instructions chosen vary on a csae-by-case basis due to the whims of the optimiser. This design in principle produces better performance than a hypothetical solution that turned PTX asm into AMD asm, because it conveniently eliminates the optimisation barrier an asm block typically represents. Care, of course, is taken to handle the wacky memory consistency concerns that this implies!

We're documenting which ones are expected to perform worse than on NVIDIA, though!

ckitching

2 वर्ष पहलेdiscuss

[I work on SCALE]

Mapping inline ptx to AMD machine code would indeed suck. Converting it to LLVM IR right at the start of compilation (when the initial IR is being generated) is much simpler, since it then gets "compiled forward" with the rest of the code. It's as if you wrote C++/intrinsics/whatever instead.

Note that nvcc accepts a different dialect of C++ from clang (and hence hipcc), so there is in fact more that separates CUDA from hip (at the language level) than just find/replace. We discuss this a little in [the manual](https://docs.scale-lang.com/manual/dialects/)

Handling differences between the atomic models is, indeed, "fun". But since CUDA is a programming language with documented semantics for its memory consistency (and so is PTX) it is entirely possible to arrange for the compiler to "play by NVIDIA's rules".