Fix UB in float-to-unsigned scalar type conversions#1283
Fix UB in float-to-unsigned scalar type conversions#1283shivansh023023 wants to merge 1 commit intoNVIDIA:mainfrom
Conversation
|
Note Reviews pausedIt looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the Use the following commands to manage reviews:
Use the checkboxes below for quick actions:
📝 WalkthroughWalkthroughAdds safe float→unsigned conversion helpers in C++ and updates Python builtin generation to emit and reference safe native names/namespaces for float→unsigned casts; expands tests to exercise uint8/uint64 float-cast edge cases. Changes
Sequence Diagram(s)sequenceDiagram
participant User as Python user code
participant BuiltinGen as WarpBuiltinGen
participant Native as C++ Native (`wp`)
participant AD as Adjoint/Runtime
User->>BuiltinGen: request scalar cast (float -> unsigned)
BuiltinGen->>BuiltinGen: detect float source & unsigned dest\ncompute safe_native (e.g., float32_to_uint8)\nset namespace/native_func accordingly
BuiltinGen->>Native: reference resolved native function
Native->>Native: perform safe conversion (NaN/bounds handling)
Native->>AD: provide no-op adjoint stub
AD->>User: gradient/adjoint behavior (no-op for these casts)
Estimated code review effort🎯 3 (Moderate) | ⏱️ ~20 minutes 🚥 Pre-merge checks | ✅ 2 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (2 passed)
✏️ Tip: You can configure your own custom pre-merge checks in the settings. ✨ Finishing Touches🧪 Generate unit tests (beta)
Comment |
Greptile SummaryThis PR fixes undefined behavior (C++ §7.3.11) when casting negative Key points from the review:
Confidence Score: 3/5
Important Files Changed
Flowchart%%{init: {'theme': 'neutral'}}%%
flowchart TD
A["wp.uint8 / uint16 / uint32 / uint64\n(float input in kernel)"] --> B{"builtins.py\ncodegen loop:\nt in unsigned_int_types\nu in float_src_types?"}
B -- "Yes" --> C["emit: float32_to_uint8(x)\n(safe_native function name,\nno wp:: namespace)"]
B -- "No" --> D["emit: wp::uint8(x)\n(direct C++ cast)"]
C --> E{"builtin.h\nsafe_float_to_int64(x)"}
E -- "NaN" --> F["return 0"]
E -- "x < -2^63" --> G["return INT64_MIN"]
E -- "x >= 2^63" --> H["return INT64_MAX"]
E -- "in range" --> I["static_cast<int64>(x)"]
F & G & H & I --> J["static_cast<uintN>(int64_val)\n✅ well-defined signed→unsigned\nmodular arithmetic"]
C --> K{"uint64 path?\nsafe_float_to_uint64(x)"}
K -- "NaN" --> L["return 0"]
K -- "x <= 0" --> M["route through\nsafe_float_to_int64(x)"]
K -- "x >= 2^64" --> N["return UINT64_MAX\n(clamp/saturate)"]
K -- "x >= 2^63" --> O["safe_float_to_int64(x - 2^63)\n+ 2^63\n✅ upper-half mapping"]
K -- "0 < x < 2^63" --> P["safe_float_to_int64(x)\ncast to uint64"]
|
![greptile-apps[bot]](https://avatars.githubusercontent.com/in/867647?s=60&v=4){kind=small}
warp/native/builtin.h
Outdated
| CUDA_CALLABLE inline uint64 float32_to_uint64(float32 x) { return static_cast<uint64>(static_cast<int64>(x)); } | ||
| CUDA_CALLABLE inline uint64 float64_to_uint64(float64 x) { return static_cast<uint64>(static_cast<int64>(x)); } |
There was a problem hiding this comment.
float_to_uint64 still triggers UB for large positive floats
The intermediate static_cast<int64>(x) is itself UB when x >= 2^63 (~9.22e18). Both float32 and float64 can represent values in the range [2^63, 2^64), which are perfectly valid uint64 values — but routing them through int64 causes the same class of undefined behaviour that this PR is trying to eliminate.
As an example, wp.uint64(1e19) would invoke float32_to_uint64(1e19f), and static_cast<int64>(1e19f) is UB because 1e19 exceeds INT64_MAX (~9.22e18).
A two-branch approach is needed:
CUDA_CALLABLE inline uint64 float32_to_uint64(float32 x) {
// Values >= 2^63 cannot pass through int64 (would be UB),
// so handle the upper half of uint64 range separately.
constexpr float32 pow2_63 = 9.223372036854776e18f;
if (x >= pow2_63) {
return static_cast<uint64>(static_cast<int64>(x - pow2_63))
+ static_cast<uint64>(9223372036854775808ULL);
}
return static_cast<uint64>(static_cast<int64>(x));
}
CUDA_CALLABLE inline uint64 float64_to_uint64(float64 x) {
constexpr float64 pow2_63 = 9.223372036854775808e18;
if (x >= pow2_63) {
return static_cast<uint64>(static_cast<int64>(x - pow2_63))
+ static_cast<uint64>(9223372036854775808ULL);
}
return static_cast<uint64>(static_cast<int64>(x));
}The uint8/uint16/uint32 conversions are safe because the maximum representable value of each (255, 65535, 4294967295) is well within INT64_MAX, so the intermediate cast to int64 is always well-defined there.
temp_test_uint8_edge_cases.py
Outdated
| import warp as wp | ||
| import numpy as np | ||
|
|
||
| wp.init() | ||
|
|
||
| @wp.kernel | ||
| def test_uint8_edge_cases( | ||
| inputs: wp.array(dtype=float), | ||
| expected: wp.array(dtype=float), | ||
| ): | ||
| tid = wp.tid() | ||
| val = inputs[tid] | ||
|
|
||
| # Cast float to uint8 | ||
| casted = wp.uint8(val) | ||
|
|
||
| # Cast back to float for assertion | ||
| out = float(casted) | ||
| wp.expect_eq(out, expected[tid]) | ||
|
|
||
| def run_test(): | ||
| # Test values: -0.1, 255.1, 128.0, -100.0 | ||
| # Expected results after float -> int64 -> uint8 conversion: | ||
| # -0.1 -> 0 -> 0 | ||
| # 255.1 -> 255 -> 255 | ||
| # 128.0 -> 128 -> 128 | ||
| # -100.0 -> -100 -> 156 (since uint8(-100) = 256 - 100 = 156) | ||
| inputs = wp.array([-0.1, 255.1, 128.0, -100.0], dtype=float) | ||
| expected = wp.array([0.0, 255.0, 128.0, 156.0], dtype=float) | ||
|
|
||
| print("Launching test_uint8_edge_cases...") | ||
| wp.launch(test_uint8_edge_cases, dim=len(inputs), inputs=[inputs, expected]) | ||
| wp.synchronize() | ||
| print("Test passed successfully!") | ||
|
|
||
| if __name__ == "__main__": | ||
| run_test() |
There was a problem hiding this comment.
Temporary debug script committed to repository root
This file appears to be a local development/debugging script and should not be committed to the repository. It duplicates logic already covered by the re-enabled test in warp/tests/test_codegen_instancing.py, and the temp_ prefix makes it clear it was only ever intended to be temporary.
Please remove this file before merging. Any edge-case values worth preserving (e.g., -0.1, 255.1, -100.0) should be added as additional wp.launch calls inside test_type_closure_scalar in the official test file.
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 2
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (1)
warp/_src/builtins.py (1)
1019-1031:⚠️ Potential issue | 🟠 MajorThe new float→unsigned path still leaves UB for NaN, infinity, and large positive values.
The helpers called at line 1019 (such as
float64_to_uint64inwarp/native/builtin.hlines 109–122) route throughstatic_cast<int64>(x)first. This cast invokes undefined behavior ifxis NaN, infinity, or outsideint64_trange. Foruint64targets, values in[2^63, 2^64)are still unsafe—the intermediateint64cast fails before the conversion touint64ever begins. Guard the helpers withstd::isfinite()checks and range validation, or avoid the signed intermediate for the positive branch.🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed. In `@warp/_src/builtins.py` around lines 1019 - 1031, The float→unsigned conversion path currently uses helpers like float64_to_uint64 (referenced by the safe_native variable in add_builtin calls) that perform a static_cast to int64_t first, which is UB for NaN/inf or values outside int64 range; fix the native helpers (e.g., float32_to_uint32/float64_to_uint64 in warp/native/builtin.h) to either (A) explicitly check std::isfinite(x) and that x is within the target unsigned range before casting and handle out-of-range/NaN/inf deterministically, or (B) avoid the signed intermediate entirely by using unsigned-range-aware logic (e.g., separate paths for x >= 0 using unsigned arithmetic), and keep builtins.py’s safe_native usage unchanged so it points to the corrected, safe helper functions.
🤖 Prompt for all review comments with AI agents
Verify each finding against the current code and only fix it if needed.
Inline comments:
In `@temp_test_uint8_edge_cases.py`:
- Around line 4-37: Move the ad-hoc script into the Warp test suite by
converting the standalone file into a unittest-style test in warp/tests (e.g.,
add a new test function in test_codegen_instancing.py or a new test_*.py),
remove wp.init(), run_test(), print() and __main__ logic, and register the
kernel test under the unittest harness; replace direct wp.launch/wp.synchronize
usages with either add_function_test(get_test_devices()) if you need
multi-device coverage or wrap kernel invocation in with wp.ScopedDevice(device):
and call wp.synchronize_device() after wp.launch, and keep the kernel name
test_uint8_edge_cases and the assertion wp.expect_eq unchanged so the same cases
are exercised.
In `@warp/native/builtin.h`:
- Around line 120-123: The four functions float16_to_uint8, float16_to_uint16,
float16_to_uint32, and float16_to_uint64 reference the float16 typedef before it
is defined; move these function definitions so they appear after the typedef
float16 (typedef half float16) or change them to operate on half instead and
then provide thin float16 wrappers after the typedef. Specifically, either
relocate the existing CUDA_CALLABLE inline float16_to_uint* functions to below
the float16 typedef, or replace their parameter type with half and call the
existing float32 conversion (e.g., float32(x)), then add wrapper functions named
float16_to_uint* that take float16 and forward to the half-based implementations
once the typedef is available.
---
Outside diff comments:
In `@warp/_src/builtins.py`:
- Around line 1019-1031: The float→unsigned conversion path currently uses
helpers like float64_to_uint64 (referenced by the safe_native variable in
add_builtin calls) that perform a static_cast to int64_t first, which is UB for
NaN/inf or values outside int64 range; fix the native helpers (e.g.,
float32_to_uint32/float64_to_uint64 in warp/native/builtin.h) to either (A)
explicitly check std::isfinite(x) and that x is within the target unsigned range
before casting and handle out-of-range/NaN/inf deterministically, or (B) avoid
the signed intermediate entirely by using unsigned-range-aware logic (e.g.,
separate paths for x >= 0 using unsigned arithmetic), and keep builtins.py’s
safe_native usage unchanged so it points to the corrected, safe helper
functions.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yml
Review profile: CHILL
Plan: Pro
Run ID: 23e8749b-b715-4228-a334-8b590b82fb7b
📒 Files selected for processing (4)
temp_test_uint8_edge_cases.pywarp/_src/builtins.pywarp/native/builtin.hwarp/tests/test_codegen_instancing.py
temp_test_uint8_edge_cases.py
Outdated
| wp.init() | ||
|
|
||
| @wp.kernel | ||
| def test_uint8_edge_cases( | ||
| inputs: wp.array(dtype=float), | ||
| expected: wp.array(dtype=float), | ||
| ): | ||
| tid = wp.tid() | ||
| val = inputs[tid] | ||
|
|
||
| # Cast float to uint8 | ||
| casted = wp.uint8(val) | ||
|
|
||
| # Cast back to float for assertion | ||
| out = float(casted) | ||
| wp.expect_eq(out, expected[tid]) | ||
|
|
||
| def run_test(): | ||
| # Test values: -0.1, 255.1, 128.0, -100.0 | ||
| # Expected results after float -> int64 -> uint8 conversion: | ||
| # -0.1 -> 0 -> 0 | ||
| # 255.1 -> 255 -> 255 | ||
| # 128.0 -> 128 -> 128 | ||
| # -100.0 -> -100 -> 156 (since uint8(-100) = 256 - 100 = 156) | ||
| inputs = wp.array([-0.1, 255.1, 128.0, -100.0], dtype=float) | ||
| expected = wp.array([0.0, 255.0, 128.0, 156.0], dtype=float) | ||
|
|
||
| print("Launching test_uint8_edge_cases...") | ||
| wp.launch(test_uint8_edge_cases, dim=len(inputs), inputs=[inputs, expected]) | ||
| wp.synchronize() | ||
| print("Test passed successfully!") | ||
|
|
||
| if __name__ == "__main__": | ||
| run_test() |
There was a problem hiding this comment.
Please fold this into the real Warp test suite instead of a standalone temp script.
The edge cases are useful, but this file is structured as an ad-hoc runner (wp.init(), run_test(), print(), __main__) rather than a normal Warp test, so it only exercises the default device and is easy for CI to miss. I'd move these cases into warp/tests/test_codegen_instancing.py or another warp/tests/test_*.py module and, if you still need an explicit flush for wp.expect_eq, do it under with wp.ScopedDevice(device): via wp.synchronize_device().
♻️ Suggested direction
- wp.init()
-
- `@wp.kernel`
- def test_uint8_edge_cases(
+ `@wp.kernel`
+ def uint8_edge_cases_kernel(
inputs: wp.array(dtype=float),
expected: wp.array(dtype=float),
):
tid = wp.tid()
val = inputs[tid]
- casted = wp.uint8(val)
- out = float(casted)
- wp.expect_eq(out, expected[tid])
+ wp.expect_eq(float(wp.uint8(val)), expected[tid])
- def run_test():
- inputs = wp.array([-0.1, 255.1, 128.0, -100.0], dtype=float)
- expected = wp.array([0.0, 255.0, 128.0, 156.0], dtype=float)
- print("Launching test_uint8_edge_cases...")
- wp.launch(test_uint8_edge_cases, dim=len(inputs), inputs=[inputs, expected])
- wp.synchronize()
- print("Test passed successfully!")
-
- if __name__ == "__main__":
- run_test()
+ def test_uint8_edge_cases(test, device):
+ with wp.ScopedDevice(device):
+ inputs = wp.array([-0.1, 255.1, 128.0, -100.0], dtype=float, device=device)
+ expected = wp.array([0.0, 255.0, 128.0, 156.0], dtype=float, device=device)
+ wp.launch(uint8_edge_cases_kernel, dim=len(inputs), inputs=[inputs, expected])
+ wp.synchronize_device()Then register it with the existing Warp unittest harness instead of keeping it as a temporary script.
As per coding guidelines: warp/tests/test_*.py: "Use unittest, not pytest, for all test files in Warp" and "Use add_function_test() only when tests need to run across multiple devices via get_test_devices()"; **/*.py: "Use wp.synchronize_device() inside wp.ScopedDevice() contexts, not wp.synchronize()."
🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed.
In `@temp_test_uint8_edge_cases.py` around lines 4 - 37, Move the ad-hoc script
into the Warp test suite by converting the standalone file into a unittest-style
test in warp/tests (e.g., add a new test function in test_codegen_instancing.py
or a new test_*.py), remove wp.init(), run_test(), print() and __main__ logic,
and register the kernel test under the unittest harness; replace direct
wp.launch/wp.synchronize usages with either
add_function_test(get_test_devices()) if you need multi-device coverage or wrap
kernel invocation in with wp.ScopedDevice(device): and call
wp.synchronize_device() after wp.launch, and keep the kernel name
test_uint8_edge_cases and the assertion wp.expect_eq unchanged so the same cases
are exercised.
There was a problem hiding this comment.
Actionable comments posted: 1
♻️ Duplicate comments (1)
warp/native/builtin.h (1)
135-138:⚠️ Potential issue | 🔴 Critical
float16_to_uint*()is declared beforefloat16exists.
typedef half float16;is only introduced at Line 215, so these inline definitions use an undeclared type and won't compile as written. Move them below the alias, or takehalfhere and add thinfloat16wrappers after the typedef.#!/bin/bash # Expect the four float16_to_uint* definitions to appear before `typedef half float16;`. rg -n '\bfloat16_to_uint(8|16|32|64)\b|\btypedef\s+half\s+float16\b|\bstruct\s+half\b' warp/native/builtin.h🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed. In `@warp/native/builtin.h` around lines 135 - 138, The inline functions float16_to_uint8, float16_to_uint16, float16_to_uint32, and float16_to_uint64 are declared before the typedef half float16 exists, causing an undeclared-type compile error; move these four CUDA_CALLABLE inline definitions so they appear after the typedef half float16; alternatively, change their parameter type to half here and then add thin wrappers named float16_to_uint* after the typedef to preserve the float16 API (refer to the symbols float16_to_uint8/16/32/64 and typedef half float16 to locate the correct places).
🤖 Prompt for all review comments with AI agents
Verify each finding against the current code and only fix it if needed.
Inline comments:
In `@warp/native/builtin.h`:
- Around line 110-131: All of the float->unsigned helpers (float32_to_uint8,
float64_to_uint8, float32_to_uint16, float64_to_uint16, float32_to_uint32,
float64_to_uint32, float32_to_uint64, float64_to_uint64) perform a
static_cast<int64>(x) which is UB for NaN/±inf and for values outside int64
range; add explicit checks before that cast: use isfinite/isnan/isinf to handle
non-finite inputs, and check numeric bounds (for small unsigned types compare
against 0 and UINTn_MAX, for uint64 use the pow2_63 branch plus INT64_MAX
limits) and clamp or return a defined saturated value instead of casting when
out-of-range so the signed cast is never applied to invalid floats.
---
Duplicate comments:
In `@warp/native/builtin.h`:
- Around line 135-138: The inline functions float16_to_uint8, float16_to_uint16,
float16_to_uint32, and float16_to_uint64 are declared before the typedef half
float16 exists, causing an undeclared-type compile error; move these four
CUDA_CALLABLE inline definitions so they appear after the typedef half float16;
alternatively, change their parameter type to half here and then add thin
wrappers named float16_to_uint* after the typedef to preserve the float16 API
(refer to the symbols float16_to_uint8/16/32/64 and typedef half float16 to
locate the correct places).
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yml
Review profile: CHILL
Plan: Pro
Run ID: 7da35086-d352-4b4c-b389-2c5d7a759e28
📒 Files selected for processing (2)
warp/native/builtin.hwarp/tests/test_codegen_instancing.py
There was a problem hiding this comment.
♻️ Duplicate comments (1)
warp/native/builtin.h (1)
140-144:⚠️ Potential issue | 🔴 Critical
float16type is used before it is defined — compilation error.These functions reference
float16on lines 141-144, but the typedeftypedef half float16;appears on line 221. This will cause a compilation error sincefloat16is not yet declared at this point.🔧 Suggested fix: Move these functions after line 221 or use `half` directly
Option 1: Move these functions after line 221 (after
typedef half float16;)Option 2: Use
halfdirectly here:-// Handle float16 source -CUDA_CALLABLE inline uint8 float16_to_uint8 (float16 x) { return float32_to_uint8 (float32(x)); } -CUDA_CALLABLE inline uint16 float16_to_uint16(float16 x) { return float32_to_uint16(float32(x)); } -CUDA_CALLABLE inline uint32 float16_to_uint32(float16 x) { return float32_to_uint32(float32(x)); } -CUDA_CALLABLE inline uint64 float16_to_uint64(float16 x) { return float32_to_uint64(float32(x)); }Then add after line 221:
// Handle float16 source (after typedef) CUDA_CALLABLE inline uint8 float16_to_uint8 (float16 x) { return float32_to_uint8 (float32(x)); } CUDA_CALLABLE inline uint16 float16_to_uint16(float16 x) { return float32_to_uint16(float32(x)); } CUDA_CALLABLE inline uint32 float16_to_uint32(float16 x) { return float32_to_uint32(float32(x)); } CUDA_CALLABLE inline uint64 float16_to_uint64(float16 x) { return float32_to_uint64(float32(x)); }🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed. In `@warp/native/builtin.h` around lines 140 - 144, The four conversion functions float16_to_uint8, float16_to_uint16, float16_to_uint32, and float16_to_uint64 are declared before the typedef typedef half float16 and thus use an unknown type; fix by either moving these four functions below the typedef (after the float16 typedef) or change their parameter type to half (e.g., use half x) here and keep the typedef placement—update the signatures consistently so the compiler sees a defined type.
🤖 Prompt for all review comments with AI agents
Verify each finding against the current code and only fix it if needed.
Duplicate comments:
In `@warp/native/builtin.h`:
- Around line 140-144: The four conversion functions float16_to_uint8,
float16_to_uint16, float16_to_uint32, and float16_to_uint64 are declared before
the typedef typedef half float16 and thus use an unknown type; fix by either
moving these four functions below the typedef (after the float16 typedef) or
change their parameter type to half (e.g., use half x) here and keep the typedef
placement—update the signatures consistently so the compiler sees a defined
type.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yml
Review profile: CHILL
Plan: Pro
Run ID: 041eeb89-46a3-40dc-8bd0-4ee9a33d78ad
📒 Files selected for processing (1)
warp/native/builtin.h
There was a problem hiding this comment.
Actionable comments posted: 1
🤖 Prompt for all review comments with AI agents
Verify each finding against the current code and only fix it if needed.
Inline comments:
In `@warp/native/builtin.h`:
- Around line 111-128: Replace the hard-coded decimal upper bound in
safe_float_to_int64 with a type-aware predecessor of 2^63: compute F pow2_63 =
static_cast<F>(9223372036854775808.0) and set max_int64 =
std::nextafter(pow2_63, static_cast<F>(-INFINITY)) (or std::nextafter(pow2_63,
static_cast<F>(0)) if -INFINITY isn't available) so the upper clamp is the
largest representable F less than 2^63; adjust safe_float_to_uint64 to reuse
this pow2_63 logic (keep the existing branches but use the computed pow2_63 and
max_uint64 derived for type F) and add regression tests for both float32 and
float64 at the 2^63 boundary to assert the correct conversions (referencing
safe_float_to_int64 and safe_float_to_uint64).
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yml
Review profile: CHILL
Plan: Pro
Run ID: edf31c49-0661-498e-bb74-4af722d52439
📒 Files selected for processing (1)
warp/native/builtin.h
There was a problem hiding this comment.
🧹 Nitpick comments (1)
warp/tests/test_codegen_instancing.py (1)
1103-1103: Minor:testparameter is unused.The static analysis flagged that the
testparameter is unused since this function useswp.expect_eqinside kernels rather thantest.assertEqual(). This is consistent with other tests in the file (e.g.,test_type_closure_vector) that follow the same pattern. No action required as this is part of theadd_function_testconvention.🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed. In `@warp/tests/test_codegen_instancing.py` at line 1103, The `test` parameter in the function test_type_closure_scalar is unused but intentionally present due to the add_function_test convention; leave the parameter as-is and do not remove it—ensure test_type_closure_scalar continues using wp.expect_eq inside kernels rather than test.assertEqual so it matches the pattern used by other tests like test_type_closure_vector.
🤖 Prompt for all review comments with AI agents
Verify each finding against the current code and only fix it if needed.
Nitpick comments:
In `@warp/tests/test_codegen_instancing.py`:
- Line 1103: The `test` parameter in the function test_type_closure_scalar is
unused but intentionally present due to the add_function_test convention; leave
the parameter as-is and do not remove it—ensure test_type_closure_scalar
continues using wp.expect_eq inside kernels rather than test.assertEqual so it
matches the pattern used by other tests like test_type_closure_vector.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yml
Review profile: CHILL
Plan: Pro
Run ID: aeb0535b-2e47-4029-81ce-280f2820e4b3
📒 Files selected for processing (2)
warp/native/builtin.hwarp/tests/test_codegen_instancing.py
shivansh023023
force-pushed
the
fix/uint8-cast-ub
branch
from
c1f91c8 to
e4b78c7
Compare
Signed-off-by: shivansh023023 <singhshivansh023@gmail.com>
shivansh023023
force-pushed
the
fix/uint8-cast-ub
branch
from
e4b78c7 to
019ebf3
Compare
1 participant
Description
This PR fixes a bug causing undefined behavior when casting negative float values to unsigned integer scalar types (like
wp.uint8) inside Warp kernels.When
wp.uint8(float_val)is used, the codegen emitted a direct C-style functional castwp::uint8(float_val). Since uint8 maps directly touint8_t(unsigned char), casting a float with a negative truncated value to this unsigned type triggers undefined behavior per the C++ standard (§7.3.11).float* -> uint*pairs during codegen, the emitted C++ code now routes via custom safe-cast functional overloads in builtin.h (e.g., float32_to_uint8). These explicitly perform a two-step cast:float -> int64 -> unsigned, leveraging C++'s well-defined modular arithmetic rule for signed-to-unsigned conversion.test_codegen_instancing.py:1099.Closes FIXME in
warp/tests/test_codegen_instancing.py:1099.Checklist
Test plan
python3 -m pytest warp/tests/test_codegen_instancing.py::TestCodeGenInstancing::test_type_closure_scalar_cpu -vpasses.wp.uint8(-0.1)yields0wp.uint8(255.1)yields255wp.uint8(128.0)yields128wp.uint8(-100.0)yields156(expected modular wrap-around)Bug fix