Reject non-None kernel return annotations

CHANGELOG.md

@@ -151,6 +151,7 @@
   ([GH-1466](https://github.com/NVIDIA/warp/issues/1466)).
 - Fix tile byte-offset overflow for arrays larger than 2 GiB
   ([GH-1422](https://github.com/NVIDIA/warp/issues/1422)).
+- Raise an error when a kernel is defined with a non-`None` return annotation ([GH-1471](https://github.com/NVIDIA/warp/issues/1471)).
 
 ### Documentation
 

warp/_src/context.py

@@ -2288,6 +2288,12 @@ def build_kernel(self, kernel):
         if kernel.adj.return_var is not None:
             raise WarpCodegenTypeError(f"'{kernel.key}': Error, kernels can't have return values")
 
+        if kernel.adj.arg_types.get("return") is not None:
+            raise WarpCodegenTypeError(
+                f"'{kernel.key}': Error, kernels can't have return values; "
+                "write results to output arguments and either omit the return annotation or use '-> None'"
+            )
+
     def build_function(self, func):
         if func in self.functions:
             return

warp/tests/test_codegen.py

@@ -565,13 +565,16 @@ def f3(x: float):
     with test.assertRaisesRegex(wp.WarpCodegenTypeError, r".*Error, kernels can't have return values"):
         wp.launch(f3, dim=1, inputs=[3.0], device=device)
 
-    # TODO: specifying a return type without returning a value is benign, but should be reported to avoid confusion
-    # @wp.kernel
-    # def f4(x: float) -> float:
-    #     return
+    @wp.kernel(module="unique")
+    def f4(x: float) -> float:
+        return
 
-    # with test.assertRaisesRegex(wp.WarpCodegenTypeError, r".*Error, kernels can't have return values"):
-    #     wp.launch(f4, dim=1, inputs=[3.0], device=device)
+    with test.assertRaisesRegex(
+        wp.WarpCodegenTypeError,
+        r".*Error, kernels can't have return values; "
+        r"write results to output arguments and either omit the return annotation or use '-> None'",
+    ):
+        wp.launch(f4, dim=1, inputs=[3.0], device=device)
 
 
 def test_error_mutating_constant_in_dynamic_loop(test, device):