Repository navigation
Integer floor division (//) truncates toward zero instead of flooring #1918
Description
Bug Description
Integer // in a Warp kernel rounds toward zero rather than toward negative infinity, so it returns the wrong value whenever exactly one operand is negative. There is no error or warning.
ast.FloorDiv maps to the floordiv built-in (warp/_src/codegen.py), which is documented as "Divide two scalars using floor division" (warp/_src/builtins.py). The floating-point overloads implement that as floorf(a / b) (warp/native/builtin.h), but the integer overload declared in DECLARE_INT_OPS is plain C division, which truncates.
Reproducer
import warp as wp
import numpy as np
@wp.kernel
def floor_div(a: wp.array(dtype=wp.int32), b: wp.array(dtype=wp.int32), out: wp.array(dtype=wp.int32)):
i = wp.tid()
out[i] = a[i] // b[i]
a = wp.array([-7, 7, -1], dtype=wp.int32)
b = wp.array([2, -2, 2], dtype=wp.int32)
out = wp.zeros(3, dtype=wp.int32)
wp.launch(floor_div, dim=3, inputs=[a, b], outputs=[out])
print(out.numpy()) # [-3 -3 0]
print(np.floor_divide(a.numpy(), b.numpy())) # [-4 -4 -1] <- expected| expression | Warp integer | Python / NumPy / Warp float |
|---|---|---|
-7 // 2 |
-3 |
-4 |
7 // -2 |
-3 |
-4 |
-1 // 2 |
0 |
-1 |
The same kernel with wp.float32 returns the correct -4.0, so integer and floating-point // disagree within Warp itself.
Impact
Any kernel computing something like index // stride, a centered-stencil offset, or a grid bin where the dividend can go negative silently gets a value that is one too high.
Existing coverage misses this because randvals() in warp/tests/test_arithmetic.py only generates positive integers (rng.integers(1, high=5)), so the existing wp.floordiv assertions never see a negative operand.
Note that mod appears to be intentionally different -- it is documented as "Modulo operation using truncated division" -- which suggests the truncation in floordiv is an oversight rather than a deliberate choice.
I have a patch and a regression test ready (verified failing before the change and passing after, across int8/int16/int32/int64) and will open a PR against this issue.
System Information
Warp 1.18.0.dev2, built from source at commit d4de134 with python build_lib.py --no-cuda.
Python 3.10.12, Ubuntu 22.04 (x86_64), GCC 11.4.0, NumPy 2.2.6.
Reproduced on the cpu device. The affected code path is shared with CUDA,
since DECLARE_INT_OPS is CUDA_CALLABLE.
Activity
Floor integer division toward negative infinity [NVIDIAGH-1918]
Floor integer division toward negative infinity [NVIDIAGH-1918]
Floor integer division toward negative infinity [NVIDIAGH-1918]
shi-eric commentedon Sep 6, 2026
This is addressed by 7784b3e.
The change updates signed integer floor division so non-exact results with opposite-sign operands round toward negative infinity, matching Python, NumPy, and Warp's floating-point floor division. Unsigned integer division is unchanged, and
wp.mod()continues to use its documented truncated-division semantics.The supporting follow-up 09d3800 is required to preserve the texture benchmarks' intended workload. Correct signed floor division introduces remainder and sign checks, but the benchmarks'
wp.tid()-derived coordinates are always nonnegative. Using unsigned coordinate arithmetic selects the simpler unsigned division path so the added correctness logic does not inflate the benchmarks or obscure texture-sampling performance.Public API behavior
Python scope
wp.array//operations on signed integer arrays now use floor semantics for negative operands.Kernel scope
//expressions andwp.floordiv()calls now round toward negative infinity on CPU and CUDA.Test coverage
warp/tests/test_arithmetic.py: newtest_floordiv_negative()coverage is registered asTestArithmetic.test_floordiv_negative_{int8,int16,int32,int64}on each available device and checks mixed signs, exact and non-exact division, zero dividends, and all four signed integer widths.query_2d()andquery_3d()inasv/benchmarks/texture.pywithout adding unit tests, which is reasonable for an ASV benchmark-only adjustment.Behavioral verification
The issue's mixed-sign reproducer now matches NumPy on CPU and CUDA. Additional checks confirmed kernel
//, kernelwp.floordiv(), and Python-side Warp array//behavior across all signed integer widths.Optimize runtime signed floor division [NVIDIAGH-1918]
Optimize nonnegative index division (NVIDIAGH-1918)
Preserve BSR and FEM performance [NVIDIAGH-1918]
Floor integer division toward negative infinity [NVIDIAGH-1918]
Optimize runtime signed floor division [NVIDIAGH-1918]