skills.rest
LuisaGroup
LuisaGroupOFFICIAL
Explore all skills by this author →
WebsiteGitHub

lc_optimize

Optimize LuisaCompute DSL kernels using warp collectives, shared memory, and command batching.

1.0k|108|Updated Nov 20, 2020
Search Tags:#shared memory#DSL#GPU optimization#luisacompute#command batching#warp collectives#kernel reduction
One-click install
npx skills add https://github.com/LuisaGroup/LuisaCompute --skill lc-optimize
Or copy as Structured Prompt for Agent▼
Please help me install this Agent Skill.
Skill: lc_optimize
Source: https://github.com/LuisaGroup/LuisaCompute/tree/main/.agents/skills/lc_optimize
Command: npx skills add https://github.com/LuisaGroup/LuisaCompute --skill lc-optimize

SYSTEM DOCUMENTATION & REQUIREMENTS

What problem does it solve?

LuisaCompute DSL kernels often bottleneck on contended atomics, block-wide reductions, and inter-thread communication, and naive implementations waste cycles on shared-memory barriers and per-command driver submission overhead.

Core Features & Use Cases

  • Warp/Wave Collective Optimization: Replace shared-memory reductions and atomics with single-instruction warp primitives like warp_active_sum, warp_prefix_sum, and warp_read_lane.
  • Shared-Memory Aggregation: Apply block-local atomic privatization, two-level warp-to-block reductions, and shared staging patterns to cut global atomic traffic from O(block_size) to O(1) per block.
  • Host-Side Command Batching: Batch stream submissions with CommandList and replace blocking synchronize() calls with async callbacks to reduce driver overhead.
  • Use Case: A softmax or stream-compaction kernel that serializes on global atomics can be rewritten so each warp reduces locally, one lane per warp updates shared memory, and a single global atomic commits the block result.

Quick Start

Optimize my LuisaCompute kernel that uses shared-memory atomics for a block-wide reduction by applying warp collectives and shared-memory privatization.

Frequently Asked Questions about lc_optimize

High-intent search queries and answers about installing and using this skill.

FAQPage Schema
How do I optimize a LuisaCompute kernel bottlenecked on atomics?▼

Replace block-level shared or global atomics with warp_active_sum per warp, have the first active lane of each warp write partials to shared memory, then issue one global atomic per block. This drops global atomic traffic from O(active threads) to O(1) per block.

How to perform a block-wide reduction in LuisaCompute DSL?▼

Use a two-level reduction: reduce within each warp using warp_active_sum with no barrier, write one partial per warp to a small Shared<T> array, call sync_block(), then reduce the few partials with a second warp collective.

When should I use shared memory instead of warp collectives in LuisaCompute?▼

Use Shared<T> when cooperation spans multiple warps, when privatizing a contended global atomic, when staging global data for reuse, or when you need arbitrary cross-thread indexing. Warp collectives are faster for single-warp reductions since they need no barrier.

Does LuisaCompute warp reduction work on all GPU backends?▼

Yes, warp collectives map to CUDA warps, HIP wavefronts, Vulkan subgroups, DirectX waves, and Metal SIMD groups. Query device.compute_warp_size() on the host and warp_lane_count() on the device instead of hardcoding 32.

Why is my LuisaCompute shared-memory reduction producing wrong results?▼

Shared memory is not self-synchronizing across warps. You need sync_block() after filling shared memory and between the read and write-back phases of each tree-reduction step, computing into a register first to avoid read-after-write hazards.

How do I reduce stream submission overhead in LuisaCompute host code?▼

Batch dependent commands into a single CommandList and commit once instead of issuing many stream << calls. Use add_callback to process results asynchronously after GPU completion instead of blocking with synchronize().

Related Skills

Explore More Software Engineering Skills→
xchang1121xchang1121
6

triton-ascend-grid-config

Optimize grid and block configurations for Triton-Ascend kernels.

Community
Advanced
mindspore-aimindspore-ai
258

tilelang-cuda-optimization

Optimize TileLang CUDA kernels and diagnose compilation and runtime issues.

Official
Advanced
xchang1121xchang1121
6

cuda-c-optimization

Optimize CUDA C kernels for performance, memory access, and numerical stability.

Community
Advanced
tensormuxtensormux
54

Port CUDA Kernel to HIP

Port CUDA kernels to HIP for ROCm-enabled AMD GPUs.

Official
Advanced