Yea this is all interesting a language research.... but that AmdGPU.jl driver is the most interesting thing about this by a good margin. CUDA lockin sucks, sucks, sucks.
Julia GPU
11–20 of 28 posts
Re: Julia GPU
#12Earlier quoted context omitted.
Most of the julia GPU stuff is being made vendor agnostic via moving infrastructure from CUDA.jl to CPUCompiler.jl. The package AMDGPU.jl is coming along very well and will be plug-able with all this stuff. https://github.com/JuliaGPU/GPUCompiler.jl https://github.com/JuliaGPU/AMDGPU.jl
So we're freed from GPU vendor lock-in by getting a programming language lock-in. Out of the frying pan, into the fire. Great.
I do not get what you are trying to say. Someone is developing an open-source solution that lets you use a vast array of different numerical accelerators more efficiently and with simpler code. All the while the work is done in the open and you can copy any subset if you want to make your own tool. Why is any of this bad?
Re: Julia GPU
#13JuliaGPU appears to be a textbook example of 'how to present a bug as a feature'. "You're a CUDA programmer. Well, why don't you learn it all over again the Julia way, and also unlearn CUDA, so the next time you have to program an nVidia GPU, you don't have a choice but to do it in Julia" Vendor lock-in FTW.
Most of the julia GPU stuff is being made vendor agnostic via moving infrastructure from CUDA.jl to CPUCompiler.jl. The package AMDGPU.jl is coming along very well and will be plug-able with all this stuff. https://github.com/JuliaGPU/GPUCompiler.jl https://github.com/JuliaGPU/AMDGPU.jl
Also, oneAPI.jl:
https://github.com/JuliaGPU/oneAPI.jl
Re: Julia GPU
#14The Nvidia proprietary bit does leave HLLs like Furhark or DSLs like Neanderthal a bit of an advantage. Are there ways of doing portable GPU things with Julia?
The Julia array abstractions make it so that most code is vendor-neutral already, and you execute on whatever GPU back-end you want by using an appropriate array type. For vendor-neutral kernel programming there's GPUArrays.jl and KernelAbstractions.jl, but both aren't currently very user friendly (but are actively used as a building block for user-facing applications and APIs).
Re: Julia GPU
#15I'm not sure I understand the bounds checking example. Can someone shed some light on how CUDA.jl does bounds checking? It's not entirely straightforward to do on a GPU.
Since we use fat array objects, and not raw pointers, we know the size of the array and can perform bounds checks at run time. We then have a mechanism to throw an exception and signal it to the CPU to display it there. That's obviously quite expensive, so you can disable it with that annotation (the Julia debug setting also controls the granularity, and thus how expensive the exception handling is). It's fairly primitive, i.e. no full-featured exception handling (for now), but has proven very useful already.
Re: Julia GPU
#16I was really confused for a moment, because the article mention CUDA a lot, which is a nvidia-specific API/framework/language. I guess that's mainly to appeal to the CUDA crowd? However, Julia being seemingly based on LLVM, interfacing it with AMD GPUs should be quite doable: > Much of the initial work focused on developing tools that make it possible to write low-level code in Julia. For example, we developed the LL…
Yes, it would be great if the .jl source code didn't even mention CUDA (but right now it does, with statements such as "using CUDA" and "CuArray(...)".)
Yes, that's fair. I focused on CUDA.jl because it is the most mature, easiest to install, etc. but as I mentioned we're actively working on generalizing that support as much as possible, and as a result support for AMD (AMDGPU.jl) and Intel (oneAPI.jl) GPUs is rapidly catching up.
Re: Julia GPU
#17I'm not sure I understand the bounds checking example. Can someone shed some light on how CUDA.jl does bounds checking? It's not entirely straightforward to do on a GPU.
Since we use fat array objects, and not raw pointers, we know the size of the array and can perform bounds checks at run time. We then have a mechanism to throw an exception and signal it to the CPU to display it there. That's obviously quite expensive, so you can disable it with that annotation (the Julia debug setting also controls the granularity, and thus how expensive the exception handling is). It's fairly prim…
How do you terminate the CUDA kernel when a bounds violation is encountered by a single thread? I don't think the CUDA API exposes a mechanism to do that safely.
Re: Julia GPU
#18Earlier quoted context omitted.
Since we use fat array objects, and not raw pointers, we know the size of the array and can perform bounds checks at run time. We then have a mechanism to throw an exception and signal it to the CPU to display it there. That's obviously quite expensive, so you can disable it with that annotation (the Julia debug setting also controls the granularity, and thus how expensive the exception handling is). It's fairly prim…
How do you terminate the CUDA kernel when a bounds violation is encountered by a single thread? I don't think the CUDA API exposes a mechanism to do that safely.
You can emit `trap` or `exit` in the PTX code (although that has exposed many bugs in the PTX assembler because it does not expect that kind of often divergent control flow). But even if you'd just have the kernel return and otherwise produce invalid results, the fact that you can then report a bounds error instead of silently corrupting data and/or generating a fatal ERROR_ILLEGAL_ACCESS (requiring an application restart) is a significant usability improvement.
Re: Julia GPU
#19> From a Python programmer perspective, how does CUDA.jl compare to PyCUDA? I think the relevant comparison today is with Numba, here's a real world recurrence analysis, @cuda.jit def _sseij(Y, I, J, O): # strides sty = cuda.blockDim.x sbx = sty * cuda.blockDim.y sby = sbx * cuda.gridDim.x sbz = sby * cuda.gridDim.y # this thread's index t = (cuda.threadIdx.x + cuda.threadIdx.y * sty + cuda.blockIdx.x * sbx + cuda.bl…
Does numbacuda allow passing buffers between kernels now? Or does it still pretty much require you to write a single superkernel in a buggy, difficult-to-debug language subset?
C isn't usually a breath of fresh air, but the last two times I tried to use nubacuda and failed (~1.5 year ago), it sure felt that way.
EDIT: yes, looks like it supports on-device buffers now. I'm still in "once bitten, twice shy" mode on account of the bugs and debug story, but I'm cautiously optimistic.
Re: Julia GPU
#20Earlier quoted context omitted.
Yes, it would be great if the .jl source code didn't even mention CUDA (but right now it does, with statements such as "using CUDA" and "CuArray(...)".)
Yes, that's fair. I focused on CUDA.jl because it is the most mature, easiest to install, etc. but as I mentioned we're actively working on generalizing that support as much as possible, and as a result support for AMD (AMDGPU.jl) and Intel (oneAPI.jl) GPUs is rapidly catching up.
This is a complete novice, ill informed, question. So forgive it in advanced, but why have an AMD specific backend at all? Couldn't you just use AMD's HIP/HIP-IFY tool on the CUDA backend and get an AMD friendly version out?
https://github.com/ROCm-Developer-Tools/HIP
I realize these sort of tools aren't magic and whatever it spites out will need work, but it seems like a really good thin starting place for AMD support with a lower overhead for growth.
After the original CUDA bits can ""cross-compile"", the workflow is greatly reduced, right?
Workflow:
- update CUDA code
- push through the HIPIFY tool
- Fix what is broken (if you can fix it on the CUDA side)
After enough iterations, the CUDA code will grow friendly to HIPification...