Fix/sycl create tasks barrier - #3
Draft
abagusetty wants to merge 6 commits into
Draft
Conversation
The for(active_y)/if(threadIdx_y==active_y) pattern in every _fill_* function was a workaround for SYCL workgroup-barrier divergence that occurs when KERNEL_SETUP_AIJ kernels (gout_stride > 1) call mask_to_index while only one y-row is inside the while loop body. Fix strategy (Option 2 / row-0 only): - Restrict pair_kl computation and keep=1 decisions to threadIdx_y==0 - Pass only blockDim_x / threadIdx_x to mask_to_index (not full 2D) - All threads (all y-rows) unconditionally reach every __syncthreads() - pair_kl0 advances by blockDim_x per iteration (same logical progress) - swap[] only needs blockDim_x entries (not blockDim_x * blockDim_y) Result: - SYCL: no deadlock, all threads reach every workgroup barrier - CUDA: identical semantics to the old serialized for(active_y) loop; scan width reduced from blockDim_x*blockDim_y to blockDim_x giving fewer barriers per scan call - Zero #ifdef USE_SYCL guards remaining in algorithmic code paths; only KERNEL_SETUP() macro definition retains the ifdef (unavoidable) Affected kernels: rys_jk_2021 (2,0,2,1), rys_jk_2111 (2,1,1,1), rys_jk_2120 (2,1,2,0), rys_jk_2210 (2,2,1,0), rys_jk_3011 (3,0,1,1), rys_jk_3110 (3,1,1,0) -- all cases where gout_stride > 1.
…uplicate nbas, uniform pair_kl0 update 1. Every early-return path now executes __syncthreads() before returning so no work-item in a sub-group is left waiting at a barrier that a divergent peer already skipped (the root cause of the PVC hang). 2. Remove duplicate `uint32_t nbas = envs.nbas;` in _fill_sr_vj_tasks (compile error). 3. Move `pair_kl0 = pair_kl1` (the "skip-to-end" early-exit inside the loop) out of the non-uniform `if (threadIdx_y == 0 && ...)` branch and into the uniform `if (t_id == 0)` block at the bottom of each loop iteration, communicated via a shared flag. This avoids a non-uniform write to a loop-condition variable that caused an extra spurious iteration on SYCL.
Supersede the row-0/do_skip_flag rewrite. Keep the two valid fixes from the previous commits (barrier before early-return; divergent for(active_y) wrapper removed) but restore the original full-width parallelization so the task-build phase is correct and performant on both CUDA and SYCL/PVC. - mask_to_index: revert to full-width signature (keep, swap, threads, t_id), scanning all blockDim_x*blockDim_y lanes. Keep the standalone #ifdef USE_SYCL `auto item` so __syncthreads() (group_barrier) has `item` in scope; every _fill_* gets `item` via KERNEL_SETUP(). - _fill_*_tasks: restore pair_kl = pair_kl0 + t_id, mask_to_index over `threads`, pair_kl0 += threads, swap[threads-1], and the uniform-value pair_kl0 = pair_kl1 skip. Drop do_skip_flag entirely. Why: the row-0-only scheme made all blockDim_y rows write swap[threadIdx_x], racing row 0's real keep with zeros (nondeterministic compaction, UB on CUDA, corrupt on PVC), and screened only blockDim_x pairs/iter -> 4x more loop iterations/barriers for gout_stride>1 kernels (e.g. rys_jk_2021). The full-width scheme is inherently barrier-uniform once the divergent wrapper is gone: the while-loop condition depends only on shared pair_kl0/ntasks, so all threads execute identical iterations and reach every __syncthreads() together. Distinct t_id lanes write distinct swap slots -> no race. Matches upstream CUDA algorithm byte-for-byte -> no numerical change, no regression. Shared-memory sizing unaffected: swap needs `threads` (<=256) ints, buffer holds thousands of ints (case 261 buflen=4736+iprim*jprim doubles).
- gint: per-file prefix on gint_kernel_L to avoid cross-file ODR collision - gvhf-rys: apply head OFFSET once inside kernel; pass bare head at all launch sites (fixes CUDA double-offset) - gvhf: add launch.cuh macros, auto-named SYCL kernels, convert int3c2e pass1/pass2 - dpnp_helper: _GPUMethodProxy wraps only routines, not callable objects
Raw libgsycl kernels (lib/*/*.cu under USE_SYCL) submit fire-and-forget on the singleton in-order queue and only receive borrowed USM pointers, so they cannot keep their argument buffers alive. dpctl frees device USM eagerly on GC (synchronous sycl::free, not queue-ordered), so a still- pending kernel can read freed memory -> intermittent GPU page fault (NotPresent/Read), e.g. sum_ejk_int3c2e_ip1 in df/grad. Add a shim layer (cupy/cuda.py) that intercepts dpnp array creation and, on GC, releases the USM via SyclQueue._submit_keep_args_alive() enqueued on the in-order master queue. The host task runs after all previously submitted kernels, so the real sycl::free happens only once those kernels complete -- matching CUDA/CuPy stream-ordered-free semantics. Frees are batched to amortize host-task cost. No new host-side synchronization is introduced; ordering relies solely on the existing in-order queue. Also: - ejk_int3c2e_ip1.cu: drop [libgsycl DEBUG] queue-ptr fprintf. - gvhf/launch.cuh: add [[intel::kernel_args_restrict]] to GVHF launches.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
No description provided.