-
Notifications
You must be signed in to change notification settings - Fork 46
Debug ci hangs #377
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
base: master
Are you sure you want to change the base?
Debug ci hangs #377
Conversation
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## master #377 +/- ##
==========================================
+ Coverage 80.27% 80.54% +0.27%
==========================================
Files 12 12
Lines 730 730
==========================================
+ Hits 586 588 +2
+ Misses 144 142 -2 ☔ View full report in Codecov by Sentry. 🚀 New features to boost your workflow:
|
fc50390 to
a100e78
Compare
425ed4a to
6140c53
Compare
7b5af2a to
9e633c7
Compare
|
Your PR requires formatting changes to meet the project's style guidelines. Click here to view the suggested changes.diff --git a/lib/cl/event.jl b/lib/cl/event.jl
index ca1edf2..6460898 100644
--- a/lib/cl/event.jl
+++ b/lib/cl/event.jl
@@ -162,7 +162,8 @@ function enqueue_marker_with_wait_list(wait_for::Vector{AbstractEvent}; queue::C
n_wait_events = cl_uint(length(wait_for))
wait_evt_ids = [evt.id for evt in wait_for]
ret_evt = Ref{cl_event}()
- clEnqueueMarkerWithWaitList(queue, n_wait_events,
+ clEnqueueMarkerWithWaitList(
+ queue, n_wait_events,
isempty(wait_evt_ids) ? C_NULL : wait_evt_ids,
ret_evt)
@return_event ret_evt[]
@@ -172,7 +173,8 @@ function enqueue_barrier_with_wait_list(wait_for::Vector{AbstractEvent}; queue::
n_wait_events = cl_uint(length(wait_for))
wait_evt_ids = [evt.id for evt in wait_for]
ret_evt = Ref{cl_event}()
- clEnqueueBarrierWithWaitList(queue, n_wait_events,
+ clEnqueueBarrierWithWaitList(
+ queue, n_wait_events,
isempty(wait_evt_ids) ? C_NULL : wait_evt_ids,
ret_evt)
@return_event ret_evt[]
@@ -185,7 +187,7 @@ function enqueue_marker(; queue::CmdQueue = queue())
end
@deprecate enqueue_marker enqueue_marker_with_wait_list
-function enqueue_wait_for_events(wait_for::Vector{T}; queue::CmdQueue = queue()) where {T<:AbstractEvent}
+function enqueue_wait_for_events(wait_for::Vector{T}; queue::CmdQueue = queue()) where {T <: AbstractEvent}
wait_evt_ids = isempty(wait_for) ? C_NULL : [pointer(evt) for evt in wait_for]
GC.@preserve wait_for begin
clEnqueueWaitForEvents(queue, length(wait_for), wait_evt_ids)
@@ -193,7 +195,7 @@ function enqueue_wait_for_events(wait_for::Vector{T}; queue::CmdQueue = queue())
end
function enqueue_wait_for_events(wait_for::AbstractEvent; queue::CmdQueue = queue())
- enqueue_wait_for_events([wait_for]; queue)
+ return enqueue_wait_for_events([wait_for]; queue)
end
function enqueue_barrier(; queue::CmdQueue = queue())
diff --git a/lib/cl/kernel.jl b/lib/cl/kernel.jl
index a7a44b5..e27439b 100644
--- a/lib/cl/kernel.jl
+++ b/lib/cl/kernel.jl
@@ -158,7 +158,8 @@ end
function enqueue_kernel(k::Kernel, global_work_size, local_work_size=nothing;
global_work_offset=nothing, wait_on::Vector{Event}=Event[],
- rng_state=false, nargs=nothing, queue::CmdQueue = queue())
+ rng_state = false, nargs = nothing, queue::CmdQueue = queue()
+ )
max_work_dim = device().max_work_item_dims
work_dim = length(global_work_size)
if work_dim > max_work_dim
@@ -224,12 +225,13 @@ function enqueue_kernel(k::Kernel, global_work_size, local_work_size=nothing;
end
ret_event = Ref{cl_event}()
- clEnqueueNDRangeKernel(queue, k, work_dim, goffset, gsize, lsize,
+ clEnqueueNDRangeKernel(
+ queue, k, work_dim, goffset, gsize, lsize,
n_events, wait_event_ids, ret_event)
return Event(ret_event[], retain=false)
end
-function enqueue_task(k::Kernel; wait_for=nothing, queue::CmdQueue = queue())
+function enqueue_task(k::Kernel; wait_for = nothing, queue::CmdQueue = queue())
n_evts = 0
evt_ids = C_NULL
#TODO: this should be split out into its own function
@@ -252,7 +254,7 @@ function call(
k::Kernel, args...; global_size = (1,), local_size = nothing,
global_work_offset = nothing, wait_on::Vector{Event} = Event[],
indirect_memory::Vector{AbstractMemory} = AbstractMemory[],
- rng_state=false, queue::CmdQueue = queue()
+ rng_state = false, queue::CmdQueue = queue()
)
set_args!(k, args...)
if !isempty(indirect_memory)
@@ -306,7 +308,7 @@ function call(
clSetKernelExecInfo(k, CL_KERNEL_EXEC_INFO_USM_PTRS_INTEL, sizeof(usm_pointers), usm_pointers)
end
end
- enqueue_kernel(k, global_size, local_size; global_work_offset, wait_on, rng_state, nargs=length(args), queue)
+ return enqueue_kernel(k, global_size, local_size; global_work_offset, wait_on, rng_state, nargs = length(args), queue)
end
# From `julia/base/reflection.jl`, adjusted to add specialization on `t`.
diff --git a/lib/cl/memory/buffer.jl b/lib/cl/memory/buffer.jl
index b9e265a..413109b 100644
--- a/lib/cl/memory/buffer.jl
+++ b/lib/cl/memory/buffer.jl
@@ -84,13 +84,15 @@ end
# reading from buffer to host array, return an event
function enqueue_read(dst::Ptr, src::Buffer, src_off::Int, nbytes::Int;
- blocking::Bool=false, wait_for::Vector{Event}=Event[],
- queue::CmdQueue = queue())
+ blocking::Bool = false, wait_for::Vector{Event} = Event[],
+ queue::CmdQueue = queue()
+ )
n_evts = length(wait_for)
evt_ids = isempty(wait_for) ? C_NULL : [pointer(evt) for evt in wait_for]
GC.@preserve wait_for begin
ret_evt = Ref{cl_event}()
- clEnqueueReadBuffer(queue, src, blocking, src_off, nbytes, dst,
+ clEnqueueReadBuffer(
+ queue, src, blocking, src_off, nbytes, dst,
n_evts, evt_ids, ret_evt)
@return_nanny_event(ret_evt[], dst)
end
@@ -100,13 +102,15 @@ enqueue_read(dst::Ptr, src::Buffer, nbytes; kwargs...) =
# writing from host array to buffer, return an event
function enqueue_write(dst::Buffer, dst_off::Int, src::Ptr, nbytes::Int;
- blocking::Bool=false, wait_for::Vector{Event}=Event[],
- queue::CmdQueue = queue())
+ blocking::Bool = false, wait_for::Vector{Event} = Event[],
+ queue::CmdQueue = queue()
+ )
n_evts = length(wait_for)
evt_ids = isempty(wait_for) ? C_NULL : [pointer(evt) for evt in wait_for]
GC.@preserve wait_for begin
ret_evt = Ref{cl_event}()
- clEnqueueWriteBuffer(queue, dst, blocking, dst_off, nbytes, src,
+ clEnqueueWriteBuffer(
+ queue, dst, blocking, dst_off, nbytes, src,
n_evts, evt_ids, ret_evt)
@return_nanny_event(ret_evt[], dst)
end
@@ -117,13 +121,15 @@ enqueue_write(dst::Buffer, src::Ptr, nbytes; kwargs...) =
# copying between two buffers, return an event
function enqueue_copy(dst::Buffer, dst_off::Int, src::Buffer, src_off::Int,
nbytes::Int; blocking::Bool=false,
- wait_for::Vector{Event}=Event[],
- queue::CmdQueue = queue())
+ wait_for::Vector{Event} = Event[],
+ queue::CmdQueue = queue()
+ )
n_evts = length(wait_for)
evt_ids = isempty(wait_for) ? C_NULL : [pointer(evt) for evt in wait_for]
GC.@preserve wait_for begin
ret_evt = Ref{cl_event}()
- clEnqueueCopyBuffer(queue, src, dst, src_off, dst_off, nbytes,
+ clEnqueueCopyBuffer(
+ queue, src, dst, src_off, dst_off, nbytes,
n_evts, evt_ids, ret_evt)
@return_event ret_evt[]
end
@@ -133,8 +139,9 @@ enqueue_copy(dst::Buffer, src::Buffer, N; kwargs...) =
# map a buffer into the host address space, returning a pointer and an event
function enqueue_map(buf::Buffer, offset::Integer, nbytes::Int, flags=:rw;
- blocking::Bool=false, wait_for::Vector{Event}=Event[],
- queue::CmdQueue = queue())
+ blocking::Bool = false, wait_for::Vector{Event} = Event[],
+ queue::CmdQueue = queue()
+ )
flags = if flags == :rw
CL_MAP_READ | CL_MAP_WRITE
elseif flags == :r
@@ -150,7 +157,8 @@ function enqueue_map(buf::Buffer, offset::Integer, nbytes::Int, flags=:rw;
evt_ids = isempty(wait_for) ? C_NULL : [pointer(evt) for evt in wait_for]
GC.@preserve wait_for begin
status = Ref{Cint}()
- ptr = clEnqueueMapBuffer(queue, buf, blocking, flags, offset, nbytes,
+ ptr = clEnqueueMapBuffer(
+ queue, buf, blocking, flags, offset, nbytes,
n_evts, evt_ids, ret_evt, status)
if status[] != CL_SUCCESS
throw(CLError(status[]))
@@ -163,8 +171,10 @@ enqueue_map(buf::Buffer, nbytes::Int, flags=:rw; kwargs...) =
enqueue_map(buf, 0, nbytes, flags; kwargs...)
# unmap a buffer, return an event
-function enqueue_unmap(buf::Buffer, ptr::Ptr; wait_for::Vector{Event}=Event[],
- queue::CmdQueue = queue())
+function enqueue_unmap(
+ buf::Buffer, ptr::Ptr; wait_for::Vector{Event} = Event[],
+ queue::CmdQueue = queue()
+ )
n_evts = length(wait_for)
evt_ids = isempty(wait_for) ? C_NULL : [pointer(evt) for evt in wait_for]
GC.@preserve wait_for begin
@@ -176,8 +186,9 @@ end
# fill a buffer with a pattern, returning an event
function enqueue_fill(buf::Buffer, offset::Integer, pattern::T, N::Integer;
- wait_for::Vector{Event}=Event[],
- queue::CmdQueue = queue()) where {T}
+ wait_for::Vector{Event} = Event[],
+ queue::CmdQueue = queue()
+ ) where {T}
nbytes = N * sizeof(T)
nbytes_pattern = sizeof(T)
@assert nbytes_pattern > 0
@@ -185,7 +196,8 @@ function enqueue_fill(buf::Buffer, offset::Integer, pattern::T, N::Integer;
evt_ids = isempty(wait_for) ? C_NULL : [pointer(evt) for evt in wait_for]
GC.@preserve begin
ret_evt = Ref{cl_event}()
- clEnqueueFillBuffer(queue, buf, [pattern],
+ clEnqueueFillBuffer(
+ queue, buf, [pattern],
nbytes_pattern, offset, nbytes,
n_evts, evt_ids, ret_evt)
@return_event ret_evt[]
diff --git a/lib/cl/memory/usm.jl b/lib/cl/memory/usm.jl
index a182cd3..9a1b572 100644
--- a/lib/cl/memory/usm.jl
+++ b/lib/cl/memory/usm.jl
@@ -176,7 +176,7 @@ end
# fill a buffer with a pattern, returning an event
function enqueue_usm_fill(ptr::Union{Ptr, CLPtr}, pattern::T, N::Integer;
- queue::CmdQueue = queue(),
+ queue::CmdQueue = queue(),
wait_for::Vector{Event}=Event[]) where {T}
nbytes = N * sizeof(T)
pattern_size = sizeof(T)
@@ -184,7 +184,8 @@ function enqueue_usm_fill(ptr::Union{Ptr, CLPtr}, pattern::T, N::Integer;
evt_ids = isempty(wait_for) ? C_NULL : [pointer(evt) for evt in wait_for]
GC.@preserve wait_for begin
ret_evt = Ref{cl_event}()
- clEnqueueMemFillINTEL(queue, ptr, Ref(pattern),
+ clEnqueueMemFillINTEL(
+ queue, ptr, Ref(pattern),
pattern_size, nbytes,
n_evts, evt_ids, ret_evt)
@return_event ret_evt[]
diff --git a/test/intrinsics.jl b/test/intrinsics.jl
index 9efd9de..d2904ae 100644
--- a/test/intrinsics.jl
+++ b/test/intrinsics.jl
@@ -18,28 +18,28 @@ const simd_ns = [2, 3, 4, 8, 16]
@testset "intrinsics" begin
@testset "barrier" begin
-@time "Barrier Local Mem fence" @on_device barrier(OpenCL.LOCAL_MEM_FENCE)
-@time "Barrier global Mem fence" @on_device barrier(OpenCL.GLOBAL_MEM_FENCE)
-@time "Barrier both Mem fence" @on_device barrier(OpenCL.LOCAL_MEM_FENCE | OpenCL.GLOBAL_MEM_FENCE)
+ @time "Barrier Local Mem fence" @on_device barrier(OpenCL.LOCAL_MEM_FENCE)
+ @time "Barrier global Mem fence" @on_device barrier(OpenCL.GLOBAL_MEM_FENCE)
+ @time "Barrier both Mem fence" @on_device barrier(OpenCL.LOCAL_MEM_FENCE | OpenCL.GLOBAL_MEM_FENCE)
-@time "WorkGroup Barrier Local Mem fence" @on_device work_group_barrier(OpenCL.LOCAL_MEM_FENCE)
-@time "WorkGroup Barrier global Mem fence" @on_device work_group_barrier(OpenCL.GLOBAL_MEM_FENCE)
-@time "WorkGroup Barrier image Mem fence" @on_device work_group_barrier(OpenCL.IMAGE_MEM_FENCE)
+ @time "WorkGroup Barrier Local Mem fence" @on_device work_group_barrier(OpenCL.LOCAL_MEM_FENCE)
+ @time "WorkGroup Barrier global Mem fence" @on_device work_group_barrier(OpenCL.GLOBAL_MEM_FENCE)
+ @time "WorkGroup Barrier image Mem fence" @on_device work_group_barrier(OpenCL.IMAGE_MEM_FENCE)
-@time "WorkGroup Barrier L/G Mem fence" @on_device work_group_barrier(OpenCL.LOCAL_MEM_FENCE | OpenCL.GLOBAL_MEM_FENCE)
-@time "WorkGroup Barrier L/I Mem fence" @on_device work_group_barrier(OpenCL.LOCAL_MEM_FENCE | OpenCL.IMAGE_MEM_FENCE)
-@time "WorkGroup Barrier L/G/I Mem fence" @on_device work_group_barrier(OpenCL.GLOBAL_MEM_FENCE | OpenCL.LOCAL_MEM_FENCE | OpenCL.IMAGE_MEM_FENCE)
+ @time "WorkGroup Barrier L/G Mem fence" @on_device work_group_barrier(OpenCL.LOCAL_MEM_FENCE | OpenCL.GLOBAL_MEM_FENCE)
+ @time "WorkGroup Barrier L/I Mem fence" @on_device work_group_barrier(OpenCL.LOCAL_MEM_FENCE | OpenCL.IMAGE_MEM_FENCE)
+ @time "WorkGroup Barrier L/G/I Mem fence" @on_device work_group_barrier(OpenCL.GLOBAL_MEM_FENCE | OpenCL.LOCAL_MEM_FENCE | OpenCL.IMAGE_MEM_FENCE)
-@time "WorkGroup Barrier Local Mem fence, work item scope" @on_device work_group_barrier(OpenCL.LOCAL_MEM_FENCE, OpenCL.memory_scope_work_item)
-@time "WorkGroup Barrier Local Mem fence, workgroup scope" @on_device work_group_barrier(OpenCL.LOCAL_MEM_FENCE, OpenCL.memory_scope_work_group)
-@time "WorkGroup Barrier Local Mem fence, device scope" @on_device work_group_barrier(OpenCL.LOCAL_MEM_FENCE, OpenCL.memory_scope_device)
-@time "Skipped" cl.memory_backend() isa cl.SVMBackend && @on_device work_group_barrier(OpenCL.LOCAL_MEM_FENCE, OpenCL.memory_scope_all_svm_devices)
-@time "WorkGroup Barrier Local Mem fence, subgroup scope" @on_device work_group_barrier(OpenCL.LOCAL_MEM_FENCE, OpenCL.memory_scope_sub_group)
+ @time "WorkGroup Barrier Local Mem fence, work item scope" @on_device work_group_barrier(OpenCL.LOCAL_MEM_FENCE, OpenCL.memory_scope_work_item)
+ @time "WorkGroup Barrier Local Mem fence, workgroup scope" @on_device work_group_barrier(OpenCL.LOCAL_MEM_FENCE, OpenCL.memory_scope_work_group)
+ @time "WorkGroup Barrier Local Mem fence, device scope" @on_device work_group_barrier(OpenCL.LOCAL_MEM_FENCE, OpenCL.memory_scope_device)
+ @time "Skipped" cl.memory_backend() isa cl.SVMBackend && @on_device work_group_barrier(OpenCL.LOCAL_MEM_FENCE, OpenCL.memory_scope_all_svm_devices)
+ @time "WorkGroup Barrier Local Mem fence, subgroup scope" @on_device work_group_barrier(OpenCL.LOCAL_MEM_FENCE, OpenCL.memory_scope_sub_group)
end
@testset "mem_fence" begin
-@warn "mem_fence"
+ @warn "mem_fence"
@on_device mem_fence(OpenCL.LOCAL_MEM_FENCE)
@on_device mem_fence(OpenCL.GLOBAL_MEM_FENCE)
@@ -56,7 +56,7 @@ end
end
@testset "atomic_work_item_fence" begin
-@warn "atomic_work_item_fence"
+ @warn "atomic_work_item_fence"
@on_device atomic_work_item_fence(OpenCL.LOCAL_MEM_FENCE, OpenCL.memory_order_relaxed, OpenCL.memory_scope_work_item)
@on_device atomic_work_item_fence(OpenCL.GLOBAL_MEM_FENCE, OpenCL.memory_order_acquire, OpenCL.memory_scope_work_group)
@@ -68,7 +68,7 @@ cl.memory_backend() isa cl.SVMBackend && @on_device atomic_work_item_fence(OpenC
end
@testset "math" begin
-@warn "math"
+ @warn "math"
@testset "unary - $T" for T in float_types
@testset "$f" for f in [
@@ -106,7 +106,7 @@ end
hypot,
(^),
]
- @warn "binary - $T, $f"
+ @warn "binary - $T, $f"
x = rand(T)
y = rand(T)
broken = ispocl && T == Float16 && f == atan
@@ -118,7 +118,7 @@ end
@testset "$f" for f in [
fma,
]
- @warn "ternary - $T, $f"
+ @warn "ternary - $T, $f"
x = rand(T)
y = rand(T)
z = rand(T)
@@ -135,7 +135,7 @@ end
OpenCL.rint,
OpenCL.rsqrt,
]
- @warn "OpenCL-specific unary - $T, $f"
+ @warn "OpenCL-specific unary - $T, $f"
x = rand(T)
broken = ispocl && T == Float16 && !(f in [OpenCL.rint, OpenCL.rsqrt])
@@ -155,7 +155,7 @@ end
OpenCL.nextafter,
OpenCL.powr,
]
- @warn "OpenCL-specific binary - $T, $f"
+ @warn "OpenCL-specific binary - $T, $f"
x = rand(T)
y = rand(T)
@@ -170,7 +170,7 @@ end
x = rand(T)
y = rand(T)
z = rand(T)
- @warn "OpenCL-specific ternary - $T"
+ @warn "OpenCL-specific ternary - $T"
@test call_on_device(OpenCL.mad, x, y, z) ≈ x * y + z
end
diff --git a/test/setup.jl b/test/setup.jl
index 1d5062e..f2eddd8 100644
--- a/test/setup.jl
+++ b/test/setup.jl
@@ -155,13 +155,13 @@ macro on_device(ex...)
$code
return
end
- @info "Pre-queue"
- queue = cl.queue()
- @info "Post-queue $(queue)"
- @opencl queue $(kwargs...) $kernel()
- @info "Post-launch"
- cl.finish(queue)
- @info "Post-sync"
+ @info "Pre-queue"
+ queue = cl.queue()
+ @info "Post-queue $(queue)"
+ @opencl queue $(kwargs...) $kernel()
+ @info "Post-launch"
+ cl.finish(queue)
+ @info "Post-sync"
end
end)
end |
8cfb4e1 to
396afa2
Compare
396afa2 to
bad16c2
Compare
Meant to open this on my fork.
I think there may be a mix of some runs taking longer than 60 minutes and some real hangs.