Repository navigation
[BUG] Warp truncates double-precision constants #485
Description
Bug Description
Double-precision constants get truncated to single-precision unless done in a specific manner.
Example
import warp as wp
@wp.kernel
def test_precision():
pi = wp.float64(wp.PI)
wp.printf("pi (kernel) = %.16g\n", pi)
print(f"pi (python) = {wp.PI}")
wp.launch(test_precision, dim=(1,), inputs=[])This produces the output:
pi (python) = 3.141592653589793
pi (kernel) = 3.141592741012573
We can see what's going on in the generated CUDA C++ code:
extern "C" __global__ void test_precision_d1c1f7f9_cuda_kernel_forward(
wp::launch_bounds_t dim)
{
for (size_t _idx = static_cast<size_t>(blockDim.x) * static_cast<size_t>(blockIdx.x) + static_cast<size_t>(threadIdx.x);
_idx < dim.size;
_idx += static_cast<size_t>(blockDim.x) * static_cast<size_t>(gridDim.x))
{
// reset shared memory allocator
wp::tile_alloc_shared(0, true);
//---------
// primal vars
const wp::float32 var_0 = 3.141592653589793;
const wp::float32 var_1 = 3.141592653589793;
wp::float64 var_2;
const wp::str var_3 = "pi (kernel) = %.16g\n";
//---------
// forward
// def test_precision(): <L 7>
// pi = wp.float64(wp.PI) <L 8>
var_2 = wp::float64(var_1);
// wp.printf("pi (kernel) = %.16g\n", pi) <L 9>
printf(var_3, var_2);
}
}The right double-precision value is being captured in the kernel, but it's assigned to a wp::float32 first, causing the truncation.
Current Workaround
A workaround is to define the variable as a wp::float64 outside the kernel:
import warp as wp
pi_fp64 = wp.float64(wp.PI)
@wp.kernel
def test_precision():
pi = wp.float64(pi_fp64)
wp.printf("pi (kernel) = %.16g\n", pi)
print(f"pi (python) = {pi_fp64.value}")
wp.launch(test_precision, dim=(1,), inputs=[])This produces the expected output:
pi (python) = 3.141592653589793
pi (kernel) = 3.141592653589793
The generated code in this case is:
extern "C" __global__ void test_precision_ae841f27_cuda_kernel_forward(
wp::launch_bounds_t dim)
{
for (size_t _idx = static_cast<size_t>(blockDim.x) * static_cast<size_t>(blockIdx.x) + static_cast<size_t>(threadIdx.x);
_idx < dim.size;
_idx += static_cast<size_t>(blockDim.x) * static_cast<size_t>(gridDim.x))
{
// reset shared memory allocator
wp::tile_alloc_shared(0, true);
//---------
// primal vars
const wp::float64 var_0 = 3.141592653589793;
wp::float64 var_1;
const wp::str var_2 = "pi (kernel) = %.16g\n";
//---------
// forward
// def test_precision(): <L 7>
// pi = wp.float64(pi_fp64) <L 8>
var_1 = wp::float64(var_0);
// wp.printf("pi (kernel) = %.16g\n", pi) <L 9>
printf(var_2, var_1);
}
}System Information
No response
Activity
c0d1f1ed commentedon Jan 13, 2026
Eric reported that this also happens for integer literals. Specifically, code like
h = h * wp.uint64(0xFF51AFD7ED558CCD)becomesWe might be able to address this by emitting
const auto var_5 = ...instead.Fix: Move hash constants to module scope for proper Warp codegen
c0d1f1ed commentedon Jan 29, 2026
This turns out to be a more complicated than expected. Just writing stuff down to sort my thoughts:
floattype) are double-precision. Even though we havewp.float32andwp.float64to discern between single- and double-precision at the Warp level, there is no single-precision literal representation.x * 2.0wherexis single-precision and2.0is interpreted as single-precision. In contrast, CUDA would implicitly promotexto double-precision and perform a double-precision multiplication. So the standard recommends using thefsuffix where applicable.We have a few design options:
floatmeans single-precision in kernels and double-precision in Python. Only explicit use ofwp.float32andwp.float64ensures a fixed format. While consistent and tractable, it's too confusing.wp.float32to stick to single precision arithmetic. While this would be closest to CUDA behavior, it's not a realistic option due to causing terrible performance in many cases.2.0is single-precision.wp.PIwould be double-precision, and we could have a single-precision variant. The issue with this is that it's hard to tell for a given expression containing literals or constants whether it produces a single- or double-precision result.wp.float64(). Basically use the cast as the inverse of the f-suffix. This is tricky because inwp.float64(2.0 / 3.0)should the division be performed in single- or double-precision? The expectation is probably double-precision but this requires either performing constant folding as part of codegen or ensuring each literal gets stored in adoubleat the CUDA level.The last option is what NumPy does: Promotion rules for Python scalars. Essentially
floatis considered a "weak" type and it only becomes strongly typed to match the precision of the arithmetic operations. It makes a lot of sense from a developer's perspective, with minimal surprising behavior.It might be feasible to implement this using a C++
weak_floatclass which can cast to a float or double based on context...shi-eric commentedon Mar 9, 2026
The scope was reduced and the precision loss issue was addressed in 0d3e1ea, planned for the 1.12.1 release.
Add CHANGELOG entry for constant precision preservation (NVIDIAGH-485)
Merge branch 'ncapens/preserve-literal' into 'main'