Julia GPU

https://notamonadtutorial.com/julia-gpu-98a461d33e21

by cdsousa • 6 years ago

131 26 6 years ago

marmaduke• 6 years 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• 6 years 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.

MayeulC2• 6 years 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• 6 years 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(...)".)

maleadt4• 6 years 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.

fulafel• 6 years 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?

maleadt4• 6 years 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).

Athas2• 6 years 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.

maleadt4• 6 years 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.

Athas2• 6 years 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.

fizixer2• 6 years 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• 6 years 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/AMDGPU.jl

fizixer2• 6 years 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.

celrod• 6 years ago

Also, oneAPI.jl: https://github.com/JuliaGPU/oneAPI.jl

jpsamaroo2• 6 years ago

Tim has been working on making it easy to use CUDA.jl and AMDGPU.jl pretty interchangeably through GPUArrays.jl, and this approach seems to be pretty extensible to other accelerators like Intel's dGPUs. KernelAbstractions.jl will also be gaining AMDGPU.jl support soon, so it'll be easy to write generic kernels without buying into a single vendor's cards.

Julia GPU

Related Stories