Skip to content

Comments

fix(conversion): lower shared S@/S! to NVVM shared address space#39

Merged
tetsuo-cpp merged 1 commit intocanonfrom
fix-shared-memory-addrspace
Feb 20, 2026
Merged

fix(conversion): lower shared S@/S! to NVVM shared address space#39
tetsuo-cpp merged 1 commit intocanonfrom
fix-shared-memory-addrspace

Conversation

@tetsuo-cpp
Copy link
Owner

Summary

  • fix shared-memory lowering for S@/S! by using NVVM shared address-space pointers
  • replace hardcoded pointer address space with NVVM::kSharedMemorySpace in ForthToMemRef
  • tighten memory-ops conversion test to require !llvm.ptr<3> for shared load/store lowering
  • this makes generated PTX use ld.shared/st.shared for tiled matmul shared buffers, avoiding illegal memory access

Test plan

  • cmake --build build --target format
  • cmake --build build --target check-warpforth
  • Verified pipeline PTX for tiled matmul emits ld.shared.u64/st.shared.u64 for shared buffers

@tetsuo-cpp tetsuo-cpp merged commit 4769da0 into canon Feb 20, 2026
1 check passed
@tetsuo-cpp tetsuo-cpp deleted the fix-shared-memory-addrspace branch February 20, 2026 17:14
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant