Sync cache modifiers branch with main and add cache modifiers to copy function

[Copilot]

Overview

This PR merges the latest changes from main into the cache-modifiers branch and implements cache modifier support for the newly added copy() function. All changes maintain backward compatibility through optional parameters.

Changes

Merged main branch

• Integrated all recent changes from main, including the new copy() function (Implements copy function #113) and its tests
• Resolved conflicts by accepting main's changes and re-applying cache modifier enhancements

Cache Modifier Implementation

Added cache modifier support to the copy() function to match the existing get() and put() functions:

@triton.jit
def copy(
    src_ptr, dst_ptr, from_rank, to_rank, cur_rank, heap_bases, mask=None,
    load_cache_modifier=None,  # NEW: Control load cache behavior
    store_cache_modifier=None  # NEW: Control store cache behavior
):
    # ...
    data = tl.load(translated_src, mask=mask, cache_modifier=load_cache_modifier)
    tl.store(translated_dst, data, mask=mask, cache_modifier=store_cache_modifier)

Supported Cache Modifiers

Load cache modifiers:

• None or ".ca" (default) - Cache at all levels (CU, L2, LLC) with LRU policy
• ".cg" - Bypasses CU (L1) cache, streams through L2
• ".cv" - Bypasses all GPU caches, fetches from system memory

Store cache modifiers:

• None or ".wb" (default) - Write-back caching at all levels
• ".cg" - Cache Global, stored through L1 → L2 → LLC
• ".cs" - Cache Streaming, bypasses L1, streamed through L2
• ".wt" - Write-Through, bypasses L1 and L2

Testing

Added comprehensive test coverage in test_copy_cache_modifiers.py that validates all combinations of load and store cache modifiers, following the same pattern as existing cache modifier tests.

Documentation

Updated function docstrings with detailed descriptions of cache modifier behavior, following the CDNA ISA specifications for instruction-level cache control.

Backward Compatibility

All changes are fully backward compatible:

• Cache modifier parameters are optional with default values of None
• Existing code continues to work without modification
• Default behavior matches previous functionality

Example Usage

@triton.jit
def kernel(src_ptr, dst_ptr, heap_bases):
    from_rank = 1
    to_rank = 0
    
    # Use streaming store to avoid cache pollution
    iris.copy(src_ptr, dst_ptr, from_rank, to_rank, to_rank, heap_bases,
              load_cache_modifier=".ca",
              store_cache_modifier=".cs")

Original prompt

---

This section details on the original issue you should resolve

<issue_title>Sync cache modifiers branch with main</issue_title>
<issue_description>Merge main into the branch, resolve any conflicts, implement modifiers for newly added copy functionn.</issue_description>

<agent_instructions>Resolve conflicts and update the branch </agent_instructions>

Comments on the Issue (you are @copilot in this section)

Fixes #229

---

💡 You can make Copilot smarter by setting up custom instructions, customizing its development environment and configuring Model Context Protocol (MCP) servers. Learn more Copilot coding agent tips in the docs.