import functools
import triton
import triton.language as tl
from triton._filecheck import filecheck_test, run_filecheck_test, run_parser
from triton.compiler.errors import CompilationError
import pytest
from typing import NamedTuple
def doesnt_compile(kernel):
@functools.wraps(kernel)
def test_fn():
with pytest.raises(triton.CompilationError):
run_parser(kernel)
return test_fn
@triton.jit
def anchor(v):
pass
@tl.core._aggregate
class Pair:
first: tl.tensor
second: tl.tensor
def __init__(self, first, second):
self.first = first
self.second = second
@triton.jit
def get_first(self):
return self.first
def get_second(self, _semantic=None):
return self.second
@triton.jit
def unpack(self):
return self.get_first(), self.get_second()
def __getitem__(self, ind: tl.constexpr, _semantic=None):
if ind == 0:
return self.first
assert ind == 1
return self.second
def __setitem__(self, ind: tl.constexpr, value, _semantic=None):
if ind == 0:
self.first = value
assert ind == 1
self.second = value
@doesnt_compile
@triton.jit
def test_assign_attribute():
scalar = 11
pair = Pair(tl.arange(0, 4), scalar)
pair.second = 42
@doesnt_compile
@triton.jit
def test_augassign_attribute():
scalar = 11
pair = Pair(tl.arange(0, 4), scalar)
pair.second += 42
@filecheck_test
@triton.jit
def test_retrieve_item():
scalar = 11
pair = Pair(tl.arange(0, 4), scalar)
anchor(pair[1])
@doesnt_compile
@triton.jit
def test_assign_item():
scalar = 11
pair = Pair(tl.arange(0, 4), scalar)
pair[1] = 42
@doesnt_compile
@triton.jit
def test_augassign_item():
scalar = 11
pair = Pair(tl.arange(0, 4), scalar)
pair[1] += 42
@filecheck_test
@triton.jit
def test_jit_method():
scalar = 11
pair = Pair(tl.arange(0, 4), scalar)
a, b = pair.unpack()
anchor(a)
anchor(b)
@tl.core._aggregate
class TypeWithJitGetItem:
value: tl.tensor
def __init__(self, value):
self.value = value
@triton.jit
def __getitem__(self, ind):
return self.value
@filecheck_test
@triton.jit
def test_jit_getitem():
v = TypeWithJitGetItem(tl.arange(0, 4))
a = v[0]
anchor(a)
@tl.core._aggregate
class TypeWithBuiltinInitializer:
value: tl.tensor
def __init__(self, _semantic=None):
self.value = tl.arange(0, 4, _semantic=_semantic)
def modify(self, value, _semantic=None):
self.value = value
@filecheck_test
@triton.jit
def test_aggregate_initializers():
value = TypeWithBuiltinInitializer()
anchor(value)
value.modify(tl.arange(4, 8))
anchor(value)
@filecheck_test
@triton.jit
def test_aggregate_modification_in_for_loop():
value = TypeWithBuiltinInitializer()
for i in range(0, 2):
value.modify(tl.arange(4, 8))
anchor(value)
@filecheck_test
@triton.jit
def test_aggregate_modification_in_while_loop():
value = TypeWithBuiltinInitializer()
i = 0
while i < 1:
i = 1
value.modify(tl.arange(4, 8))
anchor(value)
@triton.jit
def forward(arg):
return arg
@triton.jit
def list_of_functions_constexpr(arg, fns: tl.constexpr):
for i in tl.static_range(len(fns)):
fns[i](arg)
@filecheck_test
@triton.jit
def test_list_of_functions():
list_of_functions_constexpr(tl.arange(0, 4), [anchor, forward])
@triton.jit
def accumulate(a, b):
return a + b
@filecheck_test
@triton.jit
def test_call_in_loop():
acc = 0
for i in range(10):
acc = accumulate(acc, i)
@tl.core._aggregate
class FunctionParent:
@triton.jit
def function_with_name():
pass
@triton.jit
def function_with_name():
pass
@filecheck_test
@triton.jit
def test_function_name_mangling():
function_with_name()
FunctionParent.function_with_name()
@tl.core._aggregate
class AggregateWithConstexpr:
a: tl.tensor
b: tl.constexpr
def __init__(self, a, b):
self.a = a
self.b = b
@staticmethod
def create(a):
return AggregateWithConstexpr(a, tl.constexpr(42))
@triton.jit
def modify(self, a):
self.a = a
return self
@triton.jit
def add_rhs_constexpr(agg):
_ = agg.a + agg.b
@filecheck_test
@triton.jit
def test_aggregate_with_constexpr():
agg = AggregateWithConstexpr.create(tl.arange(0, 4))
add_rhs_constexpr(agg)
@triton.constexpr_function
def constexpr_function(x):
return x + 1
@filecheck_test
@triton.jit
def test_constexpr_function_from_jit():
x: tl.constexpr = constexpr_function(7)
tl.arange(0, x)
def test_constexpr_function_from_python():
assert constexpr_function(7) == 8
@triton.jit
def swap(pair):
return pair.second, pair.first
@doesnt_compile
@triton.jit
def test_assign_tuple_attrs_kernel():
p = Pair(tl.arange(0, 4), tl.arange(4, 8))
p.first, p.second = swap(p)
@doesnt_compile
@triton.jit
def test_reassign_aggregate_with_constexpr():
agg = AggregateWithConstexpr.create(tl.arange(0, 4))
agg = agg.modify(tl.arange(4, 8))
@triton.constexpr_function
def make_shape(m, n):
return (m, n)
@triton.constexpr_function
def add_shape_dims(m, n):
return m + n
@filecheck_test
@triton.jit
def test_constexpr_getitem():
shape: tl.constexpr = make_shape(4, 8)
sum: tl.constexpr = add_shape_dims(shape[0], shape[1])
tl.arange(4, sum)
@triton.constexpr_function
def make_constexpr_closure(x):
x = tl.constexpr(x)
@triton.jit
def inner(shape: tl.constexpr):
return tl.full(shape, x, dtype=tl.int32)
return inner
@filecheck_test
@triton.jit
def test_constexpr_closure():
closure: tl.constexpr = make_constexpr_closure(42)
closure((128, 128))
@triton.constexpr_function
def make_constexpr_generator(f):
f = tl.constexpr(f)
@triton.jit
def inner(lhs):
return lhs + f(lhs.shape, lhs.dtype)
return inner
@triton.jit
def inner_function(shape: tl.constexpr, dtype: tl.constexpr):
return tl.full(shape, 42, dtype)
@filecheck_test
@triton.jit
def test_constexpr_generator():
generator: tl.constexpr = make_constexpr_generator(inner_function)
lhs = tl.arange(0, 128)
generator(lhs)
@triton.constexpr_function
def Box(T):
@tl.core._aggregate
class BoxImpl:
value: T
@triton.jit
def create(value):
return BoxImpl(value)
def __init__(self, value):
self.value = value
return BoxImpl
def test_late_bound_class_reference():
TensorBox = Box(tl.tensor)
@triton.jit
def kernel():
value = TensorBox(tl.arange(0, 4))
anchor(value)
run_filecheck_test(kernel)
@triton.jit
def recursive_reduce(x):
if x.shape[0] == 1:
return x
else:
x0, x1 = x.reshape((x.shape[0] // 2, 2)).split()
return recursive_reduce(x0) + recursive_reduce(x1)
@filecheck_test
@triton.jit
def test_specialized_recursion():
x = tl.arange(0, 16)
recursive_reduce(x)
@triton.jit
def trivial_return():
return
@filecheck_test
@triton.jit
def test_call_in_while():
i = 0
while i < 10:
if i == 5:
trivial_return()
else:
trivial_return()
def test_return_in_while():
@triton.jit
def kernel():
i = 0
while i < 10:
if i == 5:
return
i += 1
with pytest.raises(CompilationError) as e:
run_parser(kernel)
assert "Cannot have `return` statements inside `while` or `for` statements in triton" in str(e.value)
class TensorPtr(NamedTuple):
test: tl.constexpr
class TestTuple(NamedTuple):
__test__ = False
test: TensorPtr
@triton.jit
def foo(test: TestTuple):
x: tl.constexpr = tl.constexpr(1)
for i in tl.range(x):
tl.static_assert(test.test.test == 1)
def test_tuple_constexpr():
test = TestTuple(test=TensorPtr(tl.constexpr(1)))
run_parser(foo, args=(test, ))
@tl.core._aggregate
class AggregateWithConstexprFunction:
val: tl.constexpr
val_squared: tl.constexpr
def __init__(self, val):
self.val = tl.constexpr(val)
self.val_squared = tl.constexpr(self.square_val())
@triton.constexpr_function
def square_val(self):
return self.val * self.val
@filecheck_test
@triton.jit
def test_aggregate_constexpr_function():
agg = AggregateWithConstexprFunction(4)
anchor(agg.val)
anchor(agg.val_squared)
anchor(agg.square_val())