Rust-GPU / Rust-GPU/rust-cuda

cp.async ergonomics: missing address-space-aware intrinsic forces inline PTX

Open
#378 0 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

Dominant language
Rust
Stars
5.4k
Forks
249
PR merge metrics
No merged PRs in 30d

Description

Summary

Using cp.async from Rust-CUDA currently requires handwritten inline PTX plus manual shared-address conversion in user kernels.

In practice, the missing piece is an address-space-aware API/intrinsic for cp.async.cg.shared.global.

Problem

For an async copy path like:

  • global source pointer
  • shared-memory destination
  • 16-byte copy (4 x f32)

users currently end up writing something like:

unsafe fn shared_addr(ptr: *mut f32) -> u32 {
    let mut addr: u64;
    asm!(
        "cvta.to.shared.u64 {dst}, {src};",
        dst = out(reg64) addr,
        src = in(reg64) ptr,
    );
    addr as u32
}

unsafe fn cp_async4(dst_shared: u32, src_global: *const f32) {
    asm!(
        "cp.async.cg.shared.global [{dst}], [{src}], 16;",
        dst = in(reg32) dst_shared,
        src = in(reg64) src_global,
    );
}

and then precompute shared base addresses in the kernel and pass integer offsets into the helper.

If we instead pass generic/raw pointers through a helper, PTX tends to contain extra address-space conversion glue around the cp.async sites.

Why this matters

This is exactly the kind of operation where users want:

  • a small, explicit intrinsic
  • correct shared/global address-space handling
  • no repeated manual inline PTX in every project

Right now, getting good code requires low-level PTX knowledge and manual control over shared address conversion.

Requested improvement

A Rust-CUDA intrinsic/helper for cp.async.cg.shared.global (or equivalent family) that:

  • represents the destination as shared-memory address space explicitly
  • avoids forcing users to manually convert *mut T into a u32 shared address
  • maps cleanly to the PTX async-copy instructions used in modern CUDA kernels

Even a low-level unsafe API would be useful if it preserves the right address-space semantics and avoids generic-pointer friction.

Contributor guide

Open the contributing guide

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

Research direction

Start by locating the Rust-CUDA intrinsic or helper definitions and existing PTX handling for cp.async; the issue does not name specific files or tests. Done means an unsafe address-space-aware cp.async.cg.shared.global API supports the described 16-byte path without requiring users to write address conversion or repeated inline PTX.

Written by the indexing model from the issue text.

Assessment

Tech stack
rust
Domain
tooling
Issue type
Feature
Difficulty
5/5
Estimated time
Over a week
Activity status
Quiet
Clarity
Mostly clear
Newbie friendliness
38/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.