Repository navigation
[WIP] New CPU codegen, new GPU codegen and GPU offloading pass (combined) - #2620
Draft
ThrudPrimrose wants to merge 1145 commits into
Draft
ThrudPrimrose wants to merge 1145 commits into
ThrudPrimrose wants to merge 1145 commits into
Conversation
The transient placement set every transient named in a device IR set to GPU_Global, and the kernel-local register rule then made a GPU_Shared tile a per-thread register. Transients already in GPU_Shared or Register keep that storage. #2613's is_shared_container is the same test for the shared case once that branch lands.
The CUDA environment put cuda_runtime.h into the host frame whatever the backend, which a HIP build cannot include. dace.h already includes the runtime of the selected backend, so the environment declares no header again, as on main.
…GPUScheduler The stream lowering helpers move into gpu_stream_scheduling.py, and the monolithic strategy class becomes AutoSingleStreamGPUScheduler(monolithic=True), which keeps its on-device requirement and its per-state synchronization. NaiveGPUStreamScheduler reads compiler.cuda.max_concurrent_streams when it assigns streams instead of when it is built, so a pipeline built before a set_temporary honors it.
…een the GPU codegens The chiplet distribution of the legacy codegen (grid padding, block index remapping and guards) becomes module functions in cuda.py that the experimental kernel scope uses as well, and kernels without a thread-block map span threads. A pooled array released early is no longer freed a second time at deallocation in either codegen. Dynamic map range connectors reach the experimental kernel arguments through the helper the legacy launch uses, and a connector named after its container is not redefined at the launch site.
SymPy's cache shares the symbol M with npbench, which declares it 64-bit, so the test failed whenever both were collected in one process.
…dex types, warp size and shared memory planning The experimental codegen declares GPU map indices with common.gpu_map_index_types, the warp size with common.gpu_warp_size and warp ids with common.gpu_thread_id_type, which drops the gpu_index_type, cuda_warp_size and hip_warp_size config keys. It places every kernel's shared memory with gpu_shared_memory.plan_gpu_shared_memory through helpers shared with the legacy codegen, so symbolic and oversized shared arrays become dynamic shared memory, launched with their size and the opt-in beyond the static limit. LiftSharedOutOfNestedSDFG is removed: nested SDFGs at any depth and sibling nested SDFGs run correctly with the planner alone (tests/codegen/gpu_nested_shared_memory_test.py). InferGPUGridAndBlockSize enumerates kernels with gpu_helpers.gpu_kernels, and the kernel signature opens its body on the same line, as in the legacy codegen.
…y, give the dynamic-map test its own M32 symbol, and expect the experimental lowering of nested device maps
# Conflicts: # dace/codegen/targets/cuda.py # tests/codegen/gpu_index_type_test.py
… stream through host map scopes MoveArrayOutOfKernel now lifts a symbolically-sized Register/Default transient inside a kernel, which would otherwise be a variable-length array nvcc rejects, and prefixes the array's subscripts in inlined tasklet bodies as well as memlets and control flow. The GPU stream array reaches a NestedSDFG inside a host map through the map's pass-through connectors instead of an edge crossing the scope, which scope_dict refused with 'Leftover nodes in queue'.
… whole identifier The per-stream connector __dace_current_stream_0 contains the legacy name as a substring, so every stream-synchronization tasklet declared an alias it never read.
…writes (#2614) ## Problem With `compiler.cuda.implementation = "experimental"`, GT4Py `concat_where` tests over dynamic domains segfault in `CompiledSDFG.fast_call`. The legacy codegen works. In the generated host code, a copy between two `GPU_Global` transients is emitted as a host `dace::CopyNDDynamic<...>::Copy(...)` over device pointers. Root cause: - `ExperimentalCUDACodeGen.copy_memory` assumes `InsertExplicitCopies` has already lifted every copy touching GPU memory to a `CopyLibraryNode`. It hands whatever is left to the CPU codegen. - `InsertExplicitCopies` leaves a direct copy implicit when `_competing_writer` reports that another edge may write an overlapping region of the destination. An *undecidable* overlap (`subsets.intersects(...) is None`) counts as a conflict. - In the GT4Py SDFG, the destination is also written by a GPU map. The two regions are disjoint in K, but their bounds are symbolic `Max`/`Min` expressions, so `intersects` returns `None`. The copy stays implicit and ends up as a host copy. ## Fix - **`InsertExplicitCopies._replace_direct_copies`:** a direct copy that involves GPU-resident storage is lifted even when a competing writer may overlap. - A plain copy edge is emitted when its source access node is visited, so it runs before every other consumer of that source. To keep that guarantee, the lifted copy gets ordering edges (empty memlets) to those consumers. This is done by the new helper `_consumers_to_precede`. - If that ordering would create a cycle with the edges carried by `_carry_write_ordering`, the copy is left implicit as before. - Copies that involve only CPU memory are unchanged, so the npbench `vadv` behaviour `_competing_writer` protects is not affected. - **`ExperimentalCUDACodeGen.copy_memory`:** a host-side copy between two access nodes that touches GPU memory now raises `CodegenError` instead of emitting a host copy that would segfault. ## Tests All new tests run without a GPU (they only generate code). - `tests/codegen/experimental_cuda_competing_copy_test.py` (new): - A reduced reproducer: a GPU map writes `C[0:Min(K, N)]` and a device-to-device copy writes `C[Max(K, 0):N]`. The test asserts the host code has no `CopyND` and does have `cudaMemcpyAsync`. - A test that a GPU copy left implicit raises `CodegenError`. - `tests/passes/insert_explicit_copies_test.py::test_device_copy_competing_with_another_write_is_lifted_ahead_of_the_source_consumers`: the copy is lifted and ordered before the source's other consumer. All three fail without the fix and pass with it. On the GT4Py SDFG, the host copy is replaced by a `cudaMemcpyDeviceToDevice` of the same size the legacy codegen emits. Locally (macOS, no GPU) the following pass: `insert_explicit_copies_test.py` including the CPU numeric runs, `intermediate_dead_store_test.py` (derived from vadv), the experimental CUDA codegen tests, and `tests/gpu_specialization`. The only failures are tests that need nvcc. ## Open points - Nothing has been run on a GPU yet. A GH200 run of the GT4Py `concat_where` tests would confirm the fix end to end. - The ordering edges only keep the guarantee the graph already gave: "before the source's other consumers". If the competing writer does not depend on the source, its order relative to the copy was never guaranteed, even for the implicit copy. - Copies staged in or out through map entries/exits (`_lift_staging_edge`) keep the old skip on a competing writer. A host-side one touching GPU memory now raises the new error instead of segfaulting. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
…r once The lane partial was initialized by fissioning the state, which is only sound once the kernel body is nested; on a flat body it split the kernel scope (wrong result or a scope error). The partial is now set to the reduction identity inside the lane scope, ordered before the strided map, and the lane reduction folds the accumulator's value, read from the node that initialized it, exactly once.
…azard SplitStateByGPUClass now depends on InsertExplicitCopies: both were bare dependencies of the scheduler, so their order followed class hashes and a copy was classified GPU in some runs only (flaky sync counts). A GPU -> host edge now syncs only when the host block reads GPU output or writes an array the queued GPU work touches; exit visibility moves to a trailing sync after a host sink that GPU work still reaches.
…and a strided map Only the innermost nested SDFG got the symbol, so a strided map two levels down failed validation with "Missing symbols on nested SDFG: ['__tid']" (polybench trmm on extended).
…odegen-into-combo # Conflicts: # dace/codegen/targets/cpu.py # dace/codegen/targets/framecode.py # dace/config_schema.yml
…] pair; test imports use the package path The legacy generator keeps main's aligned operator new/delete. The readable generator's allocation is a plain new[], so its heap_free_stmt is a plain delete[].
…, with restrict Emit `T* __restrict__ p = new T DACE_ALIGN(64)[n];` instead of a split `T *p;` + `p = new ...;`. Legacy is unchanged BY CONSTRUCTION: the new CPUCodeGen.fused_heap_declarator hook returns None in the base class, so the split path is taken verbatim; verified byte-identical over a 24-code-object corpus. Fusion requires all three: (1) the pointer is not already declared -- otherwise declare_array has hoisted `T *p = nullptr;` to an outer scope and fusing would shadow it, leaving outer accesses reading null; (2) declaration_stream IS allocation_stream -- dispatch_allocate hands out different streams only for Persistent/External, whose pointer lives in the state struct and is allocated in __dace_init, so those correctly stay split; (3) the extent is symbolic -- a constant extent names the fixed type in the declaration and GCC then rejects the load-bearing DACE_ALIGN(64) with "alignment of array elements is greater than element size" (clang accepts it; dropping the alignment would silently de-align). Two deliberate deviations from the original request: - No pointee-const (`const T* p = alloc`) is emitted: a const_init array's initializing write goes THROUGH that pointer, so it would be "assignment of read-only location". The generator's real const mechanism skips allocation entirely and fuses `const T x = expr;` at the write site, so no `= alloc` is ever const. - may_alias arrays are fused WITHOUT restrict, matching Array.as_arg / generate_nsdfg_arguments. Promising no-alias on a deliberately-aliased array is a miscompile, not a readability win.
…const_scalar_abi) The readable generator bound a read-only scalar by const VALUE (`const T x`) where legacy binds by const REFERENCE (`const T& x`). Which of the two is faster is a BACKEND artifact, not a general rule, so hardcoding either is wrong: measured on Neoverse-V2, clang vectorizes the by_ref form of a masked sum-reduction to predicated loads plus an UNPREDICATED fadd, but the by_value form -- whose distinct local copy in the inlined IR shifts the autovectorizer's cost model -- to a MERGING-PREDICATED fadd, which issues on fewer FP pipelines (~17% slower there). Elsewhere a small by-value scalar is usually the canonical fast choice. Add compiler.cpu.const_scalar_abi = by_ref | by_value, defaulting to by_ref: the readable generator then reproduces legacy's ABI (the fast form on that target) at zero readability cost, and the other binding stays one knob away. The legacy generator ignores the flag entirely and always binds by reference, so its output is byte-identical. The binding is emitted by cpp.emit_memlet_reference for a nested-SDFG argument -- the only caller passing an authoritative is_write=False -- so the tests drive it through a nested SDFG with a read-only scalar in-connector, and assert both ABIs stay bit-exact vs legacy.
Collect the code-generator tuning flags into one dict instead of scattering them across
compiler.cpu, and move const_scalar_abi into it (compiler.cpu.codegen_params.const_scalar_abi).
compiler.cpu.implementation stays outside: it SELECTS the generator, it does not tune it.
Every key in the group must state which generator(s) it affects on an "Applies to:" line -- the
group is not experimental-only (an upcoming multi-TU knob will apply to both generators), so
applicability belongs per-key rather than in the group name.
Also correct the const_scalar_abi rationale, which shipped a refuted claim. The stated mechanism
("merging-predicated fadd issues on fewer FP pipelines on Neoverse-V2") is contradicted by Arm's
Neoverse V2 Software Optimization Guide, Table 3-25: one row covers ALL FADD forms with identical
latency/throughput/pipelines, and the uop arithmetic runs the other way (the measured-faster
sel+unpredicated form costs MORE V-pipe slots than the single merging-predicated fadd). The
observation stands -- binding mode does flip the lowering, correlating with ~17% -- but the
mechanism is OPEN (predicate-pipe dependency in the accumulate chain is the leading candidate) and
code layout is not excluded. The doc now says so, and notes the sign is not fixed at all (LLVM
#121262 is the same axis with the reference winning; Lemire measured by-reference 4.5x slower),
which is the whole reason this is a knob to sweep rather than a rule to hardcode.
…int64 The <array>_idx / <array>_size helpers hardcoded `long long`. Benchmarking the alternatives on x86-64 (Zen 4) says int64 is the right default but int32 is worth exposing, so add compiler.cpu.codegen_params.index_ctype rather than a constant. int32 stays strictly opt-in and is never selected automatically. It wraps past 2**31, and the bound is on the ELEMENT count, not the byte size, so an int8 array overflows at only 2 GiB of data -- reachable, silent, and unprovable by the generator for a symbolic shape. Nor is a win the reason to reach for it: most kernels sit inside a 0.5-6% noise floor, and the deltas that clear it disagree on sign (gcc/gather -26% for int32, clang/stencil +12% against it). It is a knob to sweep. Templating or overloading the helpers on the index type was measured and rejected: GCC's identical-code folding collapses the instantiations onto one body (`<stencil3d_v3_c32>: jmp <stencil3d_v2_c32>`), and C++ deduction picks 32-bit math straight from the generated `for (auto i = 0; ...)` call sites -- silently, with no diagnostic. An explicit knob cannot do that behind the caller's back. int64 emits `long long` (byte-identical to the previous signature) and int32 emits `int`, which is exactly what those call sites already declare, so an int32 helper converts nothing. Legacy inlines its offset math and has no index type to pick, so it ignores the key and stays byte-identical.
Two experimental_readable-only knobs under compiler.cpu.codegen_params, each defaulting to today's output so legacy and the existing readable output are unchanged: - heap_ptr_restrict (restrict|none, default restrict): whether a heap array pointer is declared `T* __restrict__ p` or `T* p`. restrict is the faster bet but genuinely bimodal (rust#82834: 46 faster / 18 slower; llvm#38583 even disabled vectorization), hence a knob. Never applied to a may_alias descriptor regardless of the key -- that would be a miscompile, not a win. - index_fn_qualifier (inline_constexpr|always_inline, default inline_constexpr): the qualifier on the generated <array>_idx / <array>_size helpers. inline_constexpr is `static DACE_HDFI constexpr` (the backend already inlines these at -O2); always_inline forces it via __attribute__((always_inline)) for the case where a helper is otherwise left out of line and the un-inlined access would block vectorization. Legacy emits neither construct, so it ignores both keys and stays byte-identical.
Two knobs on the shared map-loop emitter, so they apply to BOTH generators. Each spells the identical iteration space differently and defaults to today's exact output, keeping legacy byte-identical: - loop_index_type (auto|int64|int32, default auto): the induction variable's declared type. auto deduces from the lower bound, so the usual 0 makes it int -- 32-bit index math. Widening can let TBAA disprove aliasing and sink a store out of the loop, and it changes the no-wrap facts SCEV needs, but it is also more register pressure. Distinct from index_ctype, which types the helper bodies rather than the loop variable. - loop_bound_cmp (lt|le|ne, default lt): `i < end + 1` vs `i <= end` (same range, no `+ 1` add) vs `i != end + 1`. ne is emitted ONLY for a provably unit stride. With any other stride the counter can step OVER the bound, so `i != end + 1` would never fire and the loop would not terminate; a non-unit stride falls back to lt rather than emit an infinite loop. That guard is a legality condition, not a workaround -- the strided test hangs without it. Even when legal, ne forces the optimizer to prove termination before it may vectorize or unroll, where `<` hands SCEV a trip count outright, so ne is the pessimistic end of the axis and exists to be measured, not recommended.
ne previously fell back to lt whenever the stride was not provably 1, which left
the knob unavailable on exactly the strided loops worth measuring. Normalise the
bound instead: a naive `i != end + 1` is correct only when the stride divides the
range -- otherwise the counter steps over that bound, never compares equal, and
the loop does not terminate. Emitting
begin + int_ceil(end + 1 - begin, skip) * skip
gives the first value the counter actually lands on at or past the end, which it
is guaranteed to hit exactly, for any stride. Unit stride keeps the plain end + 1
so the common case stays readable.
Tests cover both an odd and an even extent, so the stride divides the range in one
case and not the other -- the case a naive bound hangs on.
…ueezing' into ngc-merge # Conflicts: # dace/frontend/python/newast.py # dace/frontend/python/replacements/array_creation_dace.py # dace/libraries/blas/nodes/gemm.py # tests/codegen/allocation_lifetime_test.py # tests/size_scalar_shape_promotion_test.py
…apes' into ngc-merge
Ported from the extended branch with the descriptor narrowing helpers it uses; the masked-copy tests pass arrays that own their memory.
ThrudPrimrose
added a commit
that referenced
this pull request
Oct 6, 2026
Conflicts follow the PR design and keep extended's bugfixes: states() replaces all_states() and Indices is gone throughout; the reduce expansion and stack placement take the PRs' forms; extended keeps its readable code generator and tile operations, which #2620 adopts. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Remove the yapf pre-commit hook, .style.yapf and the yapf pin in the linting group.
Add the ruff-format hook next to ruff-check. line-length and quote-style carry over the yapf settings; notebooks stay unformatted as before.
ruff honours fmt: off/on at statement level only, so the parametrize decorator uses a trailing fmt: skip.
# Conflicts: # dace/autodiff/library/library.py # dace/codegen/codegen.py # dace/codegen/codeobject.py # dace/codegen/compiler.py # dace/codegen/control_flow.py # dace/codegen/instrumentation/gpu_events.py # dace/codegen/instrumentation/gpu_tx_markers.py # dace/codegen/targets/cpp.py # dace/codegen/targets/cpu.py # dace/codegen/targets/cuda.py # dace/codegen/targets/framecode.py # dace/data/core.py # dace/data/distributed.py # dace/dtypes.py # dace/frontend/common/op_repository.py # dace/frontend/python/newast.py # dace/frontend/python/parser.py # dace/frontend/python/replacements/array_creation_dace.py # dace/frontend/python/replacements/array_manipulation.py # dace/frontend/python/replacements/filtering.py # dace/libraries/blas/blas_helpers.py # dace/libraries/blas/environments/rocblas.py # dace/libraries/blas/nodes/batched_matmul.py # dace/libraries/blas/nodes/dot.py # dace/libraries/blas/nodes/einsum.py # dace/libraries/blas/nodes/gemm.py # dace/libraries/blas/nodes/gemv.py # dace/libraries/blas/nodes/ger.py # dace/libraries/blas/nodes/matmul.py # dace/libraries/fft/nodes/fft.py # dace/libraries/linalg/nodes/cholesky.py # dace/libraries/linalg/nodes/tensordot.py # dace/libraries/linalg/nodes/transpose.py # dace/libraries/mpi/nodes/allreduce.py # dace/libraries/mpi/nodes/redistribute.py # dace/libraries/mpi/nodes/reduce.py # dace/libraries/onnx/nodes/onnx_op.py # dace/libraries/onnx/nodes/onnx_op_registry.py # dace/libraries/onnx/schema.py # dace/libraries/pblas/nodes/pgemv.py # dace/libraries/sparse/nodes/csrmm.py # dace/libraries/sparse/nodes/csrmv.py # dace/libraries/standard/helper.py # dace/libraries/standard/nodes/copy/node.py # dace/libraries/standard/nodes/fill/node.py # dace/libraries/standard/nodes/reduce.py # dace/libraries/standard/reduction_planner.py # dace/libraries/stencil/stencil.py # dace/library.py # dace/memlet.py # dace/properties.py # dace/sdfg/nodes.py # dace/sdfg/sdfg.py # dace/sdfg/state.py # dace/sdfg/utils.py # dace/subsets.py # dace/symbolic.py # dace/transformation/dataflow/copy_to_device.py # dace/transformation/dataflow/copy_to_map.py # dace/transformation/dataflow/gpu_transform.py # dace/transformation/dataflow/local_storage.py # dace/transformation/dataflow/map_expansion.py # dace/transformation/dataflow/mapreduce.py # dace/transformation/dataflow/reduce_expansion.py # dace/transformation/dataflow/stream_transient.py # dace/transformation/dataflow/streaming_memory.py # dace/transformation/dataflow/strip_mining.py # dace/transformation/dataflow/tasklet_fusion.py # dace/transformation/dataflow/tiling.py # dace/transformation/dataflow/vectorization.py # dace/transformation/dataflow/warp_tiling.py # dace/transformation/estimator/enumeration/enumerator.py # dace/transformation/estimator/enumeration/greedy_enumerator.py # dace/transformation/estimator/enumeration/scoring_enumerator.py # dace/transformation/interstate/__init__.py # dace/transformation/interstate/gpu_transform_sdfg.py # dace/transformation/interstate/loop_peeling.py # dace/transformation/interstate/loop_unroll.py # dace/transformation/pass_pipeline.py # dace/transformation/passes/analysis/analysis.py # dace/transformation/passes/constant_propagation.py # dace/transformation/passes/dead_dataflow_elimination.py # dace/transformation/passes/full_map_fusion.py # dace/transformation/passes/fusion_inline.py # dace/transformation/passes/length_one_array_scalar_conversion.py # dace/transformation/passes/loop_local_memory_reduction.py # dace/transformation/passes/pattern_matching.py # dace/transformation/passes/prune_symbols.py # dace/transformation/passes/scalar_promotion.py # dace/transformation/passes/scalar_to_symbol.py # dace/transformation/passes/simplification/control_flow_raising.py # dace/transformation/passes/simplify.py # dace/transformation/subgraph/composite.py # dace/transformation/subgraph/expansion.py # dace/transformation/subgraph/stencil_tiling.py # dace/transformation/subgraph/subgraph_fusion.py # dace/transformation/transformation.py # samples/sdfg_api/stencil_boundaries.py # tests/blockreduce_cudatest.py # tests/codegen/aligned_allocation_test.py # tests/codegen/allocation_after_size_assignment_test.py # tests/codegen/allocation_lifetime_test.py # tests/codegen/cpp_test.py # tests/codegen/cuda_mempool_test.py # tests/codegen/external_memory_test.py # tests/codegen/gpu_codegen_error_checks_test.py # tests/codegen/gpu_error_drain_test.py # tests/codegen/gpu_index_type_test.py # tests/codegen/gpu_memcpy_test.py # tests/codegen/gpu_nested_grid_barrier_test.py # tests/codegen/gpu_scalar_execution_context_test.py # tests/codegen/nested_kernel_transient_test.py # tests/codegen/view_allocation_test.py # tests/conftest.py # tests/cuda_block_test.py # tests/dynamic_tb_map_cudatest.py # tests/host_map_host_data_test.py # tests/library/copy_node_test.py # tests/library/reduce_cub_workspace_test.py # tests/library/reduce_test.py # tests/multistream_custom_cudatest.py # tests/numpy/searching_test.py # tests/parse_state_struct_test.py # tests/python_frontend/callback_autodetect_test.py # tests/python_frontend/program_cache_test.py # tests/transformations/gpu_transform_test.py # tests/transformations/nest_subgraph_test.py # tests/transformations/warp_tiling_test.py
# Conflicts: # dace/codegen/codegen.py # dace/codegen/common.py # dace/codegen/targets/cpp.py # dace/codegen/targets/cpu.py # dace/codegen/targets/cuda.py # dace/codegen/targets/framecode.py # dace/frontend/common/op_repository.py # dace/frontend/python/newast.py # dace/frontend/python/replacements/array_creation_dace.py # dace/libraries/blas/blas_helpers.py # dace/libraries/blas/nodes/gemm.py # dace/libraries/blas/nodes/matmul.py # dace/libraries/linalg/nodes/cholesky.py # dace/libraries/linalg/nodes/transpose.py # dace/libraries/sparse/nodes/csrmm.py # dace/libraries/sparse/nodes/csrmv.py # dace/libraries/standard/nodes/reduce.py # dace/library.py # dace/sdfg/state.py # dace/sdfg/utils.py # dace/subsets.py # dace/symbolic.py # dace/transformation/interstate/gpu_transform_sdfg.py # samples/codegen/tensor_cores.py # tests/codegen/allocation_lifetime_test.py # tests/codegen/cuda_mempool_test.py # tests/codegen/gpu_error_drain_test.py # tests/cuda_block_test.py # tests/dynamic_tb_map_cudatest.py # tests/host_map_host_data_test.py # tests/library/copy_node_test.py # tests/numpy/array_creation_test.py # tests/parse_state_struct_test.py # tests/python_frontend/callback_autodetect_test.py # tests/python_frontend/program_cache_test.py
A TileGroup property (THREAD default, WARP, BLOCK) on every tile node decides its lowering. THREAD is today's one-thread register tile, which the vectorizer emits. dace.tile calls create BLOCK nodes: on a CPU one core runs the tile and the node lowers exactly as THREAD; in a GPU kernel a block expansion spreads the output elements over the threads of the block (thread-strided, so consecutive threads touch consecutive elements), moves its register tiles to shared memory and leaves the barriers to the GPU shared-memory synchronization. Nodes whose lanes depend on each other run on one thread. The frontend gains load, store, fill, the elementwise binary and unary calls, fma, where, sum and mma, and the docs show a blocked GEMM on both targets. Offloading no longer renames register and shared-memory transients, which have no twin.
…ileSync; Barrier library node InsertTileSync walks each GPU kernel holding WARP or BLOCK tile nodes in program order. Every tile write leaves a token for the subset it wrote and the layout that wrote it; the first node reading an overlapping subset in another layout gets a barrier right before it, as does a node overwriting a subset another thread still reads, and a loop body is walked again from what its end leaves with the loop variable one step back. Same-layout chains need no barrier, and the slots of a circular buffer carry their own tokens. The barrier is a Barrier library node of a SyncScope (warp, block, grid) lowering per GPU backend; DefaultSharedMemorySync leaves tile expansions to the pass. A block-level node with a target ISA runs the ISA call on lanes_per_thread consecutive lanes per thread, so half2 stays one instruction per thread.
…ds, wherever expansion runs TileOp.expand runs InsertTileSync before a WARP or BLOCK node lowers and expands the barriers it inserts, so an SDFG expanded by hand gets them too; a kernel already holding an expanded tile node was synchronized then and is skipped, and a Barrier already in a kernel settles the tokens. The codegen hook goes. A block tile's mask may live in shared memory like the tile. devicelevel_block_size reaches the node holding a nested SDFG through its parent references, so a nested SDFG inside a loop region no longer raises StopIteration. Tile expansions run inside kernels.
…r tests run under them
…ite conflict runs on one thread dace.tile gains prod/max/min reductions and cumsum/cumprod/cummax/cummin scans (TileScan). A BLOCK reduction or scan in a GPU kernel uses cub::BlockReduce / BlockScan (hipcub on HIP); WARP raises. Code between the tile nodes runs on every thread, so a WCR there into shared data moves into a one-thread GPU_ThreadBlock map: a loop of atomics, once per block.
# Conflicts: # dace/codegen/targets/cpu.py # dace/libraries/blas/blas_helpers.py # tests/codegen/allocation_lifetime_test.py # tests/library/matmul_unit_dim_squeeze_test.py
… CPU bit-exact comparison runs on one thread main dropped all_states() and made states() a generator. The CPU corpus comparison holds both generators to bit-exact output, which a parallel reduction only gives on one thread: its rounding follows the thread count, so ludcmp, trisolv and trmm diverged from legacy at 2 to 8 threads.
…NestedGPUDeviceMapLowering (#2534) MoveArrayOutOfKernel keeps constant-size per-thread arrays of up to register_limit (16) elements in registers and moves larger and symbolically-sized ones out of the kernel; the GPU codegen preprocess pipeline keeps calling both passes.
…path The copy analysis read a conditional without an else as its arms alone, so the copy after it assumed the arm ran. An arm that brought a host-fresh container down and updated it on the device left the copy-back after the conditional overwriting the host's newer values with stale device ones whenever the arm was skipped: CLOUDSC on canon_gpu, zqxfg under if iphase[jm - 1] == 2 once MapToForLoop left the column statements as host loops. Each such conditional gets an empty else arm for the analysis (add_fallthrough_branches, next to separate_early_returns), removed again when it receives no copy. CopyAnalysis is told the arms' entry states: an added arm uses nothing, so it does not count as a second child when the conditional's entry locations are chosen.
…MapFusionHorizontal FullMapFusion is not needed for the maps the hybrid-state wrapping creates.
This branch has not been deployed
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
3 participants
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.
Combines the readable CPU code generator, the new GPU code generator (#2259) and the GPU offloading pass on top of main, so all three run through CI together.
It also replaces the const-init marking with a small pass that promotes literal-only transients to SDFG constants.