Skip to content

[FEATURE] Support max in _to_z3 #134

@mark14wu

Description

@mark14wu
Traceback (most recent call last):
  File "/home/hwu27/workspace/triton-viz/.venv/bin/triton-sanitizer", line 10, in <module>
    sys.exit(apply())
             ^^^^^^^
  File "/home/hwu27/workspace/triton-viz/triton_viz/wrapper.py", line 58, in apply
    runpy.run_path(script, run_name="__main__")
  File "<frozen runpy>", line 286, in run_path
  File "<frozen runpy>", line 98, in _run_module_code
  File "<frozen runpy>", line 88, in _run_code
  File "cache_transform_triton.py", line 165, in <module>
    result_gold = test_get_xine_cache()
                  ^^^^^^^^^^^^^^^^^^^^^
  File "cache_transform_triton.py", line 160, in test_get_xine_cache
    cos_output, sin_output = get_xine_cache(lengths, cos_cache, sin_cache, is_prompts=is_prompts)
                             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "cache_transform_triton.py", line 102, in get_xine_cache
    prefill_cache_kernel[grid](
  File "/home/hwu27/workspace/triton-viz/.venv/lib/python3.12/site-packages/triton/runtime/jit.py", line 390, in <lambda>
    return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)
                                   ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/hwu27/workspace/triton-viz/triton_viz/core/trace.py", line 68, in run
    ret = self.interpreter_fn.run(*args, **kwargs)
          ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/hwu27/workspace/triton-viz/.venv/lib/python3.12/site-packages/triton/runtime/interpreter.py", line 1380, in run
    return GridExecutor(fn, self.arg_names, grid)(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/hwu27/workspace/triton-viz/triton_viz/core/patch.py", line 488, in _grid_executor_call
    run_grid_loops(grid)
  File "/home/hwu27/workspace/triton-viz/triton_viz/core/patch.py", line 456, in run_grid_loops
    self.fn(**call_args)
  File "cache_transform_triton.py", line 27, in prefill_cache_kernel
    cos_cache_part = tl.load(
                     ^^^^^^^^
  File "/home/hwu27/workspace/triton-viz/.venv/lib/python3.12/site-packages/triton/runtime/interpreter.py", line 781, in <lambda>
    new_member = lambda *args, member=member, **kwargs: (member(*args, **
                                                         ^^^^^^^^^^^^^^^^
  File "/home/hwu27/workspace/triton-viz/.venv/lib/python3.12/site-packages/triton/language/core.py", line 42, in wrapper
    return fn(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^
  File "/home/hwu27/workspace/triton-viz/.venv/lib/python3.12/site-packages/triton/language/core.py", line 2150, in load
    return _semantic.load(pointer, mask, other, boundary_check, padding_option, cache_modifier, eviction_policy,
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/hwu27/workspace/triton-viz/.venv/lib/python3.12/site-packages/triton/language/semantic.py", line 1086, in load
    return self._load_legacy(ptr, mask, other, boundary_check, padding, cache, eviction, is_volatile)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/hwu27/workspace/triton-viz/.venv/lib/python3.12/site-packages/triton/language/semantic.py", line 1068, in _load_legacy
    self.builder.create_masked_load(ptr.handle, mask.handle, other.handle if other else None, cache,
  File "/home/hwu27/workspace/triton-viz/triton_viz/core/patch.py", line 215, in <lambda>
    lambda *args, **kwargs: patched_op(*args, **kwargs),
                            ^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/hwu27/workspace/triton-viz/triton_viz/core/patch.py", line 187, in __call__
    ret = self.callbacks.op_overrider(*args, **kwargs)
          ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/hwu27/workspace/triton-viz/triton_viz/clients/sanitizer/sanitizer.py", line 1671, in op_load_overrider
    self._handle_access_check(ret)
  File "/home/hwu27/workspace/triton-viz/triton_viz/clients/sanitizer/sanitizer.py", line 1521, in _handle_access_check
    z3_addr, z3_constraints = expr.eval()
                              ^^^^^^^^^^^
  File "/home/hwu27/workspace/triton-viz/triton_viz/clients/sanitizer/sanitizer.py", line 1050, in eval
    expr, constraints = self._to_z3()
                        ^^^^^^^^^^^^^
  File "/home/hwu27/workspace/triton-viz/triton_viz/clients/sanitizer/sanitizer.py", line 1164, in _to_z3
    ptr, constraints_ptr = self.ptr._to_z3()
                           ^^^^^^^^^^^^^^^^^
  File "/home/hwu27/workspace/triton-viz/triton_viz/clients/sanitizer/sanitizer.py", line 1176, in _to_z3
    ptr_z3, constraints_ptr = self.ptr._to_z3()
                              ^^^^^^^^^^^^^^^^^
  File "/home/hwu27/workspace/triton-viz/triton_viz/clients/sanitizer/sanitizer.py", line 1172, in _to_z3
    self._z3, self._constraints = self.arg._to_z3()
                                  ^^^^^^^^^^^^^^^^^
  File "/home/hwu27/workspace/triton-viz/triton_viz/clients/sanitizer/sanitizer.py", line 1177, in _to_z3
    offset_z3, constraints_offset = self.offset._to_z3()
                                    ^^^^^^^^^^^^^^^^^^^^
  File "/home/hwu27/workspace/triton-viz/triton_viz/clients/sanitizer/sanitizer.py", line 1104, in _to_z3
    lhs, constraints_lhs = self.lhs._to_z3()
                           ^^^^^^^^^^^^^^^^^
  File "/home/hwu27/workspace/triton-viz/triton_viz/clients/sanitizer/sanitizer.py", line 1105, in _to_z3
    rhs, constraints_rhs = self.rhs._to_z3()
                           ^^^^^^^^^^^^^^^^^
  File "/home/hwu27/workspace/triton-viz/triton_viz/clients/sanitizer/sanitizer.py", line 1157, in _to_z3
    raise NotImplementedError("_to_z3 of max is not implemented yet")
NotImplementedError: _to_z3 of max is not implemented yet

Metadata

Metadata

Labels

No labels
No labels

Type

No type

Projects

No projects

Milestone

No milestone

Relationships

None yet

Development

No branches or pull requests

Issue actions