Skip to content

Port cache modifiers, other, and volatile to DeviceContext and Gluon APIs #470

@mawad-amd

Description

@mawad-amd

Summary

PR #158 added cache_modifier, volatile, and other parameters to the freestanding Triton JIT API (iris.load, iris.store, iris.get, iris.put, iris.copy). The DeviceContext (ctx) and Gluon APIs need to be updated to match.

Current State

Parameter Freestanding DeviceContext (ctx) Gluon
load
other
cache_modifier
volatile
store
cache_modifier
get
other
load_cache_modifier
store_cache_modifier
put
other
load_cache_modifier
store_cache_modifier
copy
other
load_cache_modifier
store_cache_modifier

What Needs to Be Done

1. DeviceContext (iris.py class DeviceContext, ~line 1526)

Add missing parameters to all 5 methods and pass them through to the underlying tl.load()/tl.store() calls:

  • load: Add other, cache_modifier, volatile → pass to tl.load(..., other=other, cache_modifier=cache_modifier, volatile=volatile)
  • store: Add cache_modifier → pass to tl.store(..., cache_modifier=cache_modifier)
  • get: Add other, load_cache_modifier, store_cache_modifier → pass load_cache_modifier to tl.load() as cache_modifier, pass store_cache_modifier to tl.store() as cache_modifier
  • put: Add other, load_cache_modifier, store_cache_modifier → same pattern as get
  • copy: Add other, load_cache_modifier, store_cache_modifier → same pattern

2. Gluon (iris/experimental/iris_gluon.py class IrisDeviceCtx, ~line 146)

Add cache modifier parameters to all 5 methods and pass them through to the underlying gl.load()/gl.store() calls:

  • load: Add cache_modifier, volatile → pass to gl.load(..., cache_modifier=cache_modifier, volatile=volatile)
  • store: Add cache_modifier → pass to gl.store(..., cache_modifier=cache_modifier)
  • get: Add load_cache_modifier, store_cache_modifier → pass to respective gl.load()/gl.store() calls
  • put: Add load_cache_modifier, store_cache_modifier → same pattern
  • copy: Add load_cache_modifier, store_cache_modifier → same pattern

Note: Check whether gl.load() and gl.store() support cache_modifier and volatile kwargs. If gluon doesn't support these yet, skip gluon and note it in the PR.

3. Tests

  • Add unit tests for the DeviceContext API mirroring the existing freestanding tests in tests/unittests/test_{load,store,get,put,copy}_cache_modifiers.py.
  • If gluon supports cache modifiers, add corresponding gluon tests as well.

Reference

  • Freestanding API signatures to match: iris/iris.py lines 1898–2227
  • Existing cache modifier tests to mirror: tests/unittests/test_*_cache_modifiers.py
  • PR Introduce cache modifiers #158 for full context on the original implementation

Metadata

Metadata

Labels

gluonirisIris project issue

Type

No type

Projects

No projects

Milestone

No milestone

Relationships

None yet

Development

No branches or pull requests

Issue actions