[Feature Request] PDL Support

[LeiWang1999]

Required prerequisites

• I have searched the Issue Tracker that this hasn't already been reported. (comment there if it has.)

Motivation

Programmatic Dependent Launch and Synchronization

helpful for us to overlap two kernels.

Solution

Introduce T.trigger_pdl and T.pdl_sync for cudaTriggerProgrammaticLaunchCompletion and cudaGridDependencySynchronize.

Introduce pdl initialization which we can borrow from T.annotate_l2_hit_ratio

[coderabbitai]

📝 CodeRabbit Plan Mode

Generate an implementation plan and prompts that you can use with your favorite coding agent.

• Create Plan

Examples

• Example 1
• Example 2

_{You can configure auto-planning by selecting labels in the issue_enrichment configuration.}

🔗 Similar Issues

Related Issues

• [BUG] Missing __syncthreads() after AtomicAdd in generated CUDA kernel #1257
• [Feature Request] check async copy is used for auto T.Pipelined #1201
• [Help] Does TileLang currently support writing efficient kernels on the Ampere architecture #1261

🔗 Related PRs

#809 - [Feature] Add ptx_cp_async_barrier_noinc intrinsic and related functionality [merged]
#1121 - [Enhancement] Add missing fence_barrier_init primitive after mbarrier init [merged]
#1272 - [Bugfix] Fix multiple cg defination when using T.sync_grid [merged]
#1414 - [Enhancement] Introduce T.__ldg [merged]

👤 Suggested Assignees

• chengyupku
• LJC00118
• daifucong233
• the-strawhat
• LeiWang1999

---

🧪 Issue enrichment is currently in open beta.

To disable automatic issue enrichment, add the following to your .coderabbit.yaml:

issue_enrichment:
  auto_enrich:
    enabled: false

💬 Have feedback or questions? Drop into our discord or schedule a call!

[w169q169]

I can come help with this

[LeiWang1999]

Thanks @w169q169 feel free to connect if you come across any issues :)

[w169q169]

@LeiWang1999 Hey! I’ve got this PR #1494 ready for review. This feature was kinda tricky for me as a beginner, so any tips or feedback would be super helpful. Thanks!