Julia implementation #106

tom91136 · 2021-06-10T04:15:27Z

This PR adds the Julia implementation of BabelStream with the following implementations:

PlainStream.jl - Single threaded for
ThreadedStream.jl - Threaded implementation with Threads.@threads macros
DistributedStream.jl - Process based parallelism with @distributed macros
CUDAStream.jl - Direct port of BabelStream's native CUDA implementation using CUDA.jl
AMDGPUStream.jl - Direct port of BabelStream's native HIP implementation using AMDGPU.jl

See README.md for details on build and run instructions.

Performance is surprisingly good across all supported hardware platforms.
All benchmarks uses Julia 1.6.1, specific versions of each package are available in Manifest.toml

AMDGPU.jl is currently still in heavy development, although the project reports most core features are working.
At the time of this PR, there are still a few issues that makes it unsuitable for production use:

Kernel performance is inconsistent (tested on on Radeon VII w/ ROCm 3.10)
API for device selection is missing, the implementation had to resort to introspection
Rapid kernel submission seems to cause ref-count overflow issues if we attempt to wait for each kernel to complete, this is a bit counter-intuitive. The workaround is to use hardware-based events for synchronisation.

Finally, there isn't a process-based (e.g MPI) implementation of BabelStream so comparison for DistributedStream.jl has been omitted.
That said, performance seems to be significantly worst than ThreadedStream.jl due to the added serialisation overhead.

Future work

We should be able to include oneAPi.jl once it is ready for general use.

There's also OpenCL.jl but it simply wraps the OpenCL host API; kernels must still be written in OpenCL C, so this wouldn't be any different from BabelStream's OCLStream.

tomdeakin

Interesting to see this - think it's pretty close. In general I'd like to make sure this is idiomatic of Julia code and we're not introducing any extra layers of abstraction in the benchmark.

JuliaStream.jl/README.md

JuliaStream.jl/src/AMDGPUStream.jl

JuliaStream.jl/src/CUDAStream.jl

JuliaStream.jl/src/DistributedStream.jl

JuliaStream.jl/src/PlainStream.jl

JuliaStream.jl/src/Stream.jl

tomdeakin · 2021-06-28T14:32:05Z

JuliaStream.jl/src/StreamData.jl

@@ -0,0 +1,8 @@
+


Why do we need to wrap up the data arrays in a whole new struct?

It's similar to what we have for Rust, there are no classes in Julia, only functions and data (structs).
The alternative is to pass everything as parameters:

function mul!(data::StreamData{T}) where {T}

becomes

function mul!(a::C, b::C, c::C, scalar::T, size::Int) where {T,C<:AbstractArray{T}}

Same applies when I copy everything back to host:

const ROCData = StreamData{T,ROCArray{T}} where {T} ... function read_data(data::ROCData{T})::VectorData{T} where {T} return VectorData{T}(data.a, data.b, data.c, data.scalar, data.size) end

read_data would be very long, and it needs to be implemented for 5 different implementations.

I think I prefer passing things as parameters.

Show the selected implementation and not a constant "threaded"

tom91136 · 2021-06-30T18:32:46Z

Ready for review again.

tomdeakin · 2021-07-01T11:26:10Z

Thanks @tom91136. I think I prefer the parameter passing rather than making a structure just to hold the arrays. I think that in a larger code with more arrays, you're just going to have to pass things around rather than keep wrapping things up in bundles to pass to different functions. Ideally we're aiming to write BabelStream in a way that is representative of something much bigger.

giordano · 2021-08-12T12:08:31Z

JuliaStream.jl/src/CUDAStream.jl

+const DotBlocks = 256::Int
+
+function devices()
+  return !CUDA.functional(false) ? [] :


Minor nitpick: make the function type-stable

Suggested change

return !CUDA.functional(false) ? [] :

return !CUDA.functional(false) ? String[] :

giordano · 2021-08-12T12:10:41Z

JuliaStream.jl/src/AMDGPUStream.jl

+    map(repr, gpu_agents_in_repr_order())
+  catch
+    # probably unsupported
+    []


Suggested change

[]

String[]

giordano · 2021-08-12T12:13:26Z

JuliaStream.jl/src/DistributedStream.jl

+  end
+end
+
+function copy!(data::SharedArrayData{T}) where {T}


Note that here and in the functions below where T isn't really used to constrain the signature nor in the body of the function, the parameter is redundant and

Suggested change

function copy!(data::SharedArrayData{T}) where {T}

function copy!(data::SharedArrayData)

would work just fine

giordano · 2021-08-12T12:22:28Z

JuliaStream.jl/src/Stream.jl

+  if config.float
+    type = Float32
+  else
+    type = Float64
+  end


For the record, you can also do

Suggested change

if config.float

type = Float32

else

type = Float64

end

type = if config.float

Float32

else

Float64

end

or even

Suggested change

if config.float

type = Float32

else

type = Float64

end

type = config.float ? Float32 : Float64

if you want to be even more concise. Also in the if below you can save a couple of benchmark = in the same way

giordano · 2021-08-12T12:29:07Z

Performance is surprisingly good across all supported hardware platforms.

😃

Is there something we can do to move this forward? I had a very quick look, but could do a more thorough review, if that helps

tom91136 · 2021-08-12T13:03:56Z

@giordano Thanks for the review! This PR will be used for an upcoming submission to PMBS so I got a few more local changes (I've added a functional oneAPI.jl and KA implementation) that I'm in the process of finalising. I'll incorporate your review and put up a final version for further review by the end of the week.
If you're interested, the PMBS submission will also include a compute bound benchmark written in Julia.

@tomdeakin and I had a discussion on the parameter passing and I think we've settled on the current approach being acceptable.

jpsamaroo

Cool stuff! Let me know if I can help with anything AMDGPU.jl-related (I'm the author and maintainer of it).

jpsamaroo · 2021-08-12T14:24:15Z

JuliaStream.jl/src/AMDGPUStream.jl

+    return
+  end
+  AMDGPU.wait(
+    soft = false, # soft wait causes HSA_REFCOUNT overflow issues


This shouldn't be an issue on AMDGPU.jl 0.2.11, where I've disabled HSA_REFCOUNT until I can figure out why it's broken.

Still, hard wait could be more efficient than soft (since it hypothetically can use interrupts), but I'd make sure to test with both.

I've tested it with and without soft wait, and for the very few times soft=true actually run to completion (in most cases it dies after 70 iterations, so 5*70 kernels) the performance seems within margin of error. I think I'll stick to soft=false for now.

Can you explain what you mean by:

actually run to completion (in most cases it dies after 70 iterations, so 5*70 kernels)

If you have a stacktrace or other information, that would help greatly.

Isolate projects to avoid transitive dependency Add parameter for passing devices Incorporate further reviews Update all dependencies

maleadt

Looks good! I added a couple of small comments. Great to see that the performance is good 🙂

maleadt · 2021-08-17T14:41:41Z

JuliaStream.jl/src/CUDAStream.jl

+end
+
+function init_arrays!(data::CuData{T}, _, init::Tuple{T,T,T}) where {T}
+  CUDA.fill!(data.a, init[1])


You shouldn't need the module prefix here, as fill! is just a Base method that dispatches correctly since you're passing a CuArray argument already. Only when that's not the case (e.g. with fill, which just takes an integer) you need the CUDA-specific version.

maleadt · 2021-08-17T14:45:57Z

JuliaStream.jl/src/CUDAStream.jl

+
+function copy!(data::CuData{T}, _) where {T}
+  function kernel(a, c)
+    i = (blockIdx().x - 1) * blockDim().x + threadIdx().x # only blockIdx starts at 1


threadIdx also starts at 1, but our arrays are 1-indexed, so everything works out here.

maleadt · 2021-08-17T14:47:20Z

JuliaStream.jl/src/CUDAStream.jl

+    return
+  end
+  @cuda blocks = data.size ÷ TBSize threads = TBSize kernel(data.b, data.c, data.scalar)
+  CUDA.synchronize()


I'm not familiar with the underlying architecture here, but do you really need to synchronize after every kernel? That's not great for performance. Our GPU-to-CPU copy already synchronizes, so typically you don't need anything else.

The original implementation issues cudaDeviceSynchronize(); after kernel submission (see

BabelStream/CUDAStream.cu

Line 161 in 8f9ca7b

cudaDeviceSynchronize();

) so this port must match that. Does the synchronize in Julia does a lot more than cudaDeviceSynchronize()? From what I can see, synchronisation doesn't seem to add any more overhead compare to the native version.

It should have similar overhead.

maleadt · 2021-08-17T14:48:08Z

JuliaStream.jl/src/CUDAStream.jl

+  end
+  partial_sum = CuArray{T}(undef, DotBlocks)
+  @cuda blocks = DotBlocks threads = TBSize kernel(data.a, data.b, data.size, partial_sum)
+  CUDA.synchronize()


The sum below will run on the same stream as this kernel, so you don't need the synchronization.

maleadt · 2021-08-17T14:54:11Z

JuliaStream.jl/src/oneAPIStream.jl

+end
+
+function add!(data::oneData{T}, groupsize::Int) where {T}
+  function kernel(a, b, c)


It's a very minor detail, but you might want to put the T typevar on the inner kernel function instead, e.g., by using a::AbstractArray{T}, ... as arguments. That way you enforce specialization of the kernel method, whereas the outer CPU method doesn't necessarily need to be specialized (and you can use eltype(data) if you need the element type there).

But these things only matter when you work with very complicated inputs.

Does this apply to all other GPU implementations or just oneAPI?
Ideally this would be an exemplar for writing kernels so would you recommend I annotate every other GPU implementation like this?

Aside:
I thought about doing this with more precise types (e.g ::CuDeviceVector{T} for CUDA) but felt like the actual type is an implementation detail for the kernel launch macro (@cuda, @oneapi, etc) so I guess AbstractArray{T} is the best we an do.

This applies to all Julia code, not just the oneAPI GPU implementation. If you add a typevar (::T .. .where T) you force the compiler to specialize. That's needed for the GPU kernel, because we don't want or support dynamic code for GPUs, but the host function (here, add!) may remain underspecialized if the compiler thinks that's useful.

And yeah, the AbstractArray suggestion is so that you don't have to deal with the exact device-side types, which are an implementation detail.

maleadt · 2021-08-17T14:56:37Z

JuliaStream.jl/src/CUDAStream.jl

+    error("arraysize ($(arraysize)) must be divisible by $(TBSize)!")
+  end
+
+  # so CUDA's device is 0 indexed, so -1 from Julia


CUDA.jl's device! is 0-indexed too.

Oops, the comment was outdated as we switched to passing device instances instead of index around.

Drop soft=false for AMDGPU as this option was removed Update dependencies

Update dependencies

vchuravy · 2021-08-28T10:18:46Z

JuliaStream.jl/src/ThreadedStream.jl

+  # This doesn't do well on aarch64 because of the excessive Threads.threadid() ccall 
+  # and inhibited vectorisation from the lack of @fastmath 
+  # partial = zeros(T, Threads.nthreads())
+  # Threads.@threads for i = 1:data.size


Normally I would do a:

@threads for tid in 1:nthreads() region = .... acc = 0.0 @simd for i in region acc += ... end partial[i] = acc end

Or using @tkf FLoops.jl/Folds

Is that semantically equivalent to static_par_ranged? (minus the @simd bit)

I think so. Basically that's the canonical example of saying, I want one iteration per thread. Then you have to do the splitting manually, but you also don't have to reimplement the lower-level details xD

https://github.com/JuliaFolds/FLoops.jl#parallel-loop is an alternative

Yeah, it'd be better to use @threads for i 1:nthreads() than touching internals like task.sticky = true and jl_set_task_tid.

The main difference to FLoops.jl is that it lets the Julia runtime dynamic load-balancing. It doesn't matter if you are only micro-benchmarking a parallel dot product but static scheduling like @threads for and static_par_ranged would fail when used inside other parallel functions (it has to wait for all worker threads to be available). So, if the aim is to benchmark practical and composable pieces of routines in Julia, it may be better to use dynamic scheduling. On the other hand, if you want to benchmark what you can squeeze out from the Julia runtime at the moment, static scheduling like @threads for and static_par_ranged may be fine.

Another difference to FLoops.jl is that, by default, it uses divide-and-conquer (DAC) scheduling (@spawns and waits) while @threads for and static_par_ranged use sequential scheduling. I find that DAC (in Julia) sometimes helps performance since it has larger parallelism (in the sense of work/span ratio), especially when there are a lot of cores.

vchuravy · 2021-08-28T10:19:59Z

JuliaStream.jl/src/ThreadedStream.jl

+  partial = Vector{T}(undef,  Threads.nthreads())
+  static_par_ranged(data.size, Threads.nthreads()) do group, startidx, endidx
+    acc = zero(T)
+    @fastmath for i = startidx:endidx


Use @simd instead of @fastmath that makes sure that only the reduction variable get's reordered

Thanks! I didn't really want to use @fastmath in a benchmark like STREAM and I'm happy report I get the same runtime with @simd on both ThunderX2 and A64fx!

On the A64FX node of Isambard, I've got a good speed up with @simd when setting the environment variable JULIA_LLVM_ARGS="-aarch64-sve-vector-bits-min=512": JuliaLang/julia#40308 (comment). Note that this requires llvm 12, so Julia v1.7

Yeah @simd finds the reduction chain, and then adds a subset of the fast math flags onto that chain.

For the A64fx you might also want to set JULIA_LLVM_ARGS="-aarch64-sve-vector-bits-min=512" see JuliaLang/julia#40308 (comment)

BTW, I didn't show in that comment that there is also a good scaling with the size of the floating type (including half-precision Float16):

$ JULIA_LLVM_ARGS="-aarch64-sve-vector-bits-min=512" julia -q julia> using BenchmarkTools julia> function sumsimd(x) s = zero(eltype(x)) @simd for xi in x s += xi end s end sumsimd (generic function with 1 method) julia> @btime sumsimd(x) setup=(x = randn(Float64, 1_000_000)) 191.912 μs (0 allocations: 0 bytes) 1853.0335322487956 julia> @btime sumsimd(x) setup=(x = randn(Float32, 1_000_000)) 80.330 μs (0 allocations: 0 bytes) 400.9806f0 julia> @btime sumsimd(x) setup=(x = randn(Float16, 1_000_000)) 42.761 μs (0 allocations: 0 bytes) Float16(1.872e3)

Related to this, you can watch this lightning talk at JuliaCon 2021 about speedup on using Float16 on Isambard: https://www.youtube.com/watch?v=btHfZr2C0GA

Note that the speedup reported at JuliaCon requires JuliaLang/julia#40216

tom91136 added 5 commits June 10, 2021 04:20

Initial Julia implementation

b3efa6a

set pwd to JuliaStream.jl for CI run

63f471f

Use addprocs() for DistributedStream

2cf8ca5

Drop -p N for DistributedStream.jl CI

c5ad3f3

Larger arraysize for CI

d799535

tom91136 added the enhancement label Jun 10, 2021

tom91136 requested a review from tomdeakin June 10, 2021 04:15

tomdeakin added this to the v4.0 milestone Jun 22, 2021

tomdeakin requested changes Jun 28, 2021

View reviewed changes

tom91136 added 6 commits June 30, 2021 18:09

Inline AMDGPU's hard_wait

4e6c567

Show the selected implementation and not a constant "threaded"

Merge branch 'main' into julia

fe18065

Switch back to -p for DistributedStream

d675875

Use -p 2 and no arg for JuliaStream in CI

4183155

Inline blocks in CUDAStream

2e957d3

Add comment about blockIdx/workgroupIdx in Julia

7c1e04a

giordano reviewed Aug 12, 2021

View reviewed changes

jpsamaroo reviewed Aug 12, 2021

View reviewed changes

Add oneAPI and KA implementation

a26699c

Isolate projects to avoid transitive dependency Add parameter for passing devices Incorporate further reviews Update all dependencies

maleadt reviewed Aug 17, 2021

View reviewed changes

tom91136 mentioned this pull request Aug 17, 2021

Unable to build AMDGPU.jl with roc* libraries JuliaGPU/AMDGPU.jl#154

Closed

tom91136 added 5 commits August 18, 2021 01:59

Update PlainStream with context

bb271dd

Address CUDA comments

c445b64

Drop soft=false for AMDGPU as this option was removed Update dependencies

Add type annotation for all kernels

4853457

Update dependencies

Use custom static reduction for CPU

13cb8ff

Pause GC during benchmark to reduce noise

41f1767

vchuravy reviewed Aug 28, 2021

View reviewed changes

Use @simd instead of @fastmath for CPU reduction

78b52a4

giordano mentioned this pull request Nov 18, 2021

LLVM vector intrinsic support for SVE? JuliaLang/julia#40308

Open

tomdeakin merged commit d1b1722 into main Nov 25, 2021

	return !CUDA.functional(false) ? [] :
	return !CUDA.functional(false) ? String[] :

	function copy!(data::SharedArrayData{T}) where {T}
	function copy!(data::SharedArrayData)

Julia implementation #106

Julia implementation #106

Uh oh!

Conversation

tom91136 commented Jun 10, 2021

Future work

Uh oh!

tomdeakin left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

tom91136 commented Jun 30, 2021

Uh oh!

tomdeakin commented Jul 1, 2021

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

giordano commented Aug 12, 2021

Uh oh!

tom91136 commented Aug 12, 2021

Uh oh!

jpsamaroo left a comment • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

maleadt left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

jpsamaroo left a comment •

edited

Loading

tom91136 Aug 28, 2021 •

edited

Loading

tkf Aug 28, 2021 •

edited

Loading