fix: kernel return validation

CHANGELOG.md

@@ -198,6 +198,16 @@
 - Fix `wp.constant(wp.int32(IntEnum_value))` emitting the symbolic enum name instead of the
   integer value in generated C++/CUDA code on Python 3.10, causing compilation failures
   ([newton#2363](https://github.com/newton-physics/newton/issues/2363)).
+- Fix `ValueError: Cell is empty` crash during eager module hashing when a kernel closure references a variable
+  assigned later in the enclosing scope ([GH-913](https://github.com/NVIDIA/warp/issues/913)).
+- Fix wheels missing the bundled `libdevice.10.bc`, which caused all CUDA kernels compiled via the
+  Clang/LLVM toolchain (`llvm_cuda=True`) to fail when no system CUDA toolkit was available.
+- Fix a rare data corruption in peer-to-peer `wp.copy()` with non-contiguous sources
+  ([GH-1384](https://github.com/NVIDIA/warp/issues/1384)).
+- Fix unresolved Python C API symbols in `warp.so`/`warp.dll`/`libwarp.dylib`, allowing the library
+  to be loaded from non-Python C++ applications via `dlopen(..., RTLD_LAZY)` / `LoadLibrary()`
+  ([GH-1399](https://github.com/NVIDIA/warp/issues/1399)).
+- Fix Raise a `WarpCodegenTypeError` when a kernel specifies a non-`None` return type annotation but contains only a bare return.
 
 ### Documentation
 

warp/_src/context.py

@@ -2274,7 +2274,7 @@ def build_kernel(self, kernel):
 
         kernel.adj.build(self)
 
-        if kernel.adj.return_var is not None:
+        if kernel.adj.return_var is not None or "return" in kernel.adj.arg_types:
             raise WarpCodegenTypeError(f"'{kernel.key}': Error, kernels can't have return values")
 
     def build_function(self, func):

warp/tests/test_codegen.py

@@ -554,13 +554,19 @@ 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)
+
+    @wp.kernel(module="unique")
+    def f5(x: float) -> None:
+        return
 
-    # with test.assertRaisesRegex(wp.WarpCodegenTypeError, r".*Error, kernels can't have return values"):
-    #     wp.launch(f4, dim=1, inputs=[3.0], device=device)
+    # -> none should remain accepted
+    wp.launch(f5, dim=1, inputs=[3.0], device=device)
 
 
 def test_error_mutating_constant_in_dynamic_loop(test, device):