Skip to content

Commit e6bd9d3

Browse files
committed
[PRNG] Add check to PRNG to make sure that unsigned integer arithmetic is wrapping
1 parent d7a9a7c commit e6bd9d3

File tree

1 file changed

+12
-0
lines changed

1 file changed

+12
-0
lines changed

python/tvm/topi/random/kernel.py

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -234,6 +234,18 @@ def gen_ir(gen_ptr, out_gen_ptr, out_array_ptr):
234234
out_gen = irb.buffer_ptr(out_gen_ptr)
235235
out_array = irb.buffer_ptr(out_array_ptr)
236236

237+
# Check that unsigned arithmetic wraps, as it is required to implement threefry correctly.
238+
irb.emit(
239+
tvm.tir.AssertStmt(
240+
tvm.tir.const(0xFFFFFFFFFFFFFFFF, "uint64") + tvm.tir.const(1, "uint64")
241+
== tvm.tir.const(0, "uint64"),
242+
tvm.tir.StringImm(
243+
"Unsigned interger arithmetic is not wrapping, but threefry requires wrapping."
244+
),
245+
tvm.tir.Evaluate(0),
246+
)
247+
)
248+
237249
# Create a temporary array to hold the generator state we will use to create the random
238250
# numbers. We cannot use gen because we may need to update the key + path if there is not
239251
# enough room in the counter.

0 commit comments

Comments
 (0)