9 ms·

Julia GPU

131 points by cdsousa 6y ago

marmaduke 6y ago
> 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.blockIdx.y * sby + cuda.blockIdx.z * sbz) if t < I.size: i = I[t] j = J[t] x = nb.float32(0.0) for k in range(Y.shape[0]): x += (Y[k, i] - Y[k, j])**2 O[t] = math.sqrt(x) most of it is index calculation, but super easy
- jjoonathan 6y ago
  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.
MayeulC 6y ago
I 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 LLVM.jl package that gives us access to the LLVM APIs. Recently, our focus has shifted towards generalizing this functionality so that other GPU back-ends, like AMDGPU.jl or oneAPI.jl can benefit from developments to CUDA.jl. Vendor-neutral array operations, for examples, are now implemented in GPUArrays.jl whereas shared compiler functionality now lives in GPUCompiler.jl. That should make it possible to work on several GPU back-ends, even though most of them are maintained by only a single developer. This is the takeaway to me. First-class access to GPU accelerators using the same syntax, regardless of the vendor :)
- amelius 6y ago
  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(...)".)
  - maleadt 6y ago
    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.
    - soganess 6y ago
      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 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...
      - jpsamaroo 6y ago
        > 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? HIP and HIPify only work on C++ source code, via a Perl script. Since we start with plain Julia code, and we already have LLVM integrated into Julia's compiler, it's easiest to just change the LLVM "target" from Native to AMDGPU (or NVPTX in CUDA.jl's case) to get native machine code, while preserving Julia's semantics for the most part. Also, interfacing to ROCR (AMD's implementation of the Heterogeneous System Architecture or HSA runtime) was really easy when I first started on this, and codegen through Julia's compiler and LLVM is trivial when you have CUDAnative.jl (CUDA.jl's predecessor) to look at :) I should also mention that not everything that CUDA does maps well to AMD GPU; CUDA's streams are generally in-order (blocking), whereas AMD's queues are non-blocking unless barriers are scheduled. Also, things like hostcall (calling a CPU function from the GPU) doesn't have an obvious alternative with CUDA.
fulafel 6y ago
The 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?
- maleadt 6y ago
  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).
Athas 6y ago
I'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.
- maleadt 6y ago
  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.
  - Athas 6y ago
    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.
    - maleadt 6y ago
      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.
fizixer 6y ago
JuliaGPU 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.
- eigenspace 6y ago
  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/GPUCompiler.jl https://github.com/JuliaGPU/AMDGPU.jl https://github.com/JuliaGPU/AMDGPU.jl
  - fizixer 6y ago
    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.
    - dahart 6y ago
      What do you mean? What is programming language lock-in, and how is that worse than (or even comparable to) vendor lock-in? What's your goal?
    - krastanov 6y ago
      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?
    - mrfox321 6y ago
      What is programming language lock-in? When any code is written in some language?
monadic2 6y ago
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.

131 points by cdsousa 6y ago

> 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.blockIdx.y * sby + cuda.blockIdx.z * sbz) if t < I.size: i = I[t] j = J[t] x = nb.float32(0.0) for k in range(Y.shape[0]): x += (Y[k, i] - Y[k, j])**2 O[t] = math.sqrt(x) most of it is index calculation, but super easy

jjoonathan 6y ago
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.

MayeulC 6y ago

I 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 LLVM.jl package that gives us access to the LLVM APIs. Recently, our focus has shifted towards generalizing this functionality so that other GPU back-ends, like AMDGPU.jl or oneAPI.jl can benefit from developments to CUDA.jl. Vendor-neutral array operations, for examples, are now implemented in GPUArrays.jl whereas shared compiler functionality now lives in GPUCompiler.jl. That should make it possible to work on several GPU back-ends, even though most of them are maintained by only a single developer. This is the takeaway to me. First-class access to GPU accelerators using the same syntax, regardless of the vendor :)

amelius 6y ago
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(...)".)
- maleadt 6y ago
  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.
  - soganess 6y ago
    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 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...
    - jpsamaroo 6y ago
      > 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? HIP and HIPify only work on C++ source code, via a Perl script. Since we start with plain Julia code, and we already have LLVM integrated into Julia's compiler, it's easiest to just change the LLVM "target" from Native to AMDGPU (or NVPTX in CUDA.jl's case) to get native machine code, while preserving Julia's semantics for the most part. Also, interfacing to ROCR (AMD's implementation of the Heterogeneous System Architecture or HSA runtime) was really easy when I first started on this, and codegen through Julia's compiler and LLVM is trivial when you have CUDAnative.jl (CUDA.jl's predecessor) to look at :) I should also mention that not everything that CUDA does maps well to AMD GPU; CUDA's streams are generally in-order (blocking), whereas AMD's queues are non-blocking unless barriers are scheduled. Also, things like hostcall (calling a CPU function from the GPU) doesn't have an obvious alternative with CUDA.

fulafel 6y ago

The 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?

maleadt 6y ago
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).

Athas 6y ago

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

maleadt 6y ago
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.
- Athas 6y ago
  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.
  - maleadt 6y ago
    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.

fizixer 6y ago

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

eigenspace 6y ago
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/GPUCompiler.jl https://github.com/JuliaGPU/AMDGPU.jl https://github.com/JuliaGPU/AMDGPU.jl
- fizixer 6y ago
  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.
  - dahart 6y ago
    What do you mean? What is programming language lock-in, and how is that worse than (or even comparable to) vendor lock-in? What's your goal?
  - krastanov 6y ago
    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?
  - mrfox321 6y ago
    What is programming language lock-in? When any code is written in some language?

monadic2 6y ago

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.