diff --git a/CHANGELOG.md b/CHANGELOG.md index c24df4f70..7ec2abf81 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -11,9 +11,11 @@ compiled `(compiled)`, loaded from the cache `(cached)`, or was unable to be loaded `(error)`. - `wp.config.verbose = True` now also prints out a message upon the entry to a `wp.ScopedTimer`. -- Add additional documentation and examples demonstrating `wp.copy()`, `wp.clone()`, and `array.assign()` differentiability. +- Add additional documentation and examples demonstrating `wp.copy()`, `wp.clone()`, and `array.assign()` differentiability - Fix adding `__new__()` methods for all class `__del__()` methods to anticipate when a class instance is created but not instantiated before garbage collection. +- Add code-completion support for wp.config variables. +- Add `wp.clear_kernel_cache()` to the public API. This is equivalent to `wp.build.clear_kernel_cache()`. - Add code-completion support for `wp.config` variables. - Remove usage of a static task (thread) index for CPU kernels to address multithreading concerns. - The `mask` argument to `wp.sim.eval_fk` now accepts both integer and bool arrays diff --git a/docs/basics.rst b/docs/basics.rst index 1c815b537..2b042edfe 100644 --- a/docs/basics.rst +++ b/docs/basics.rst @@ -205,6 +205,7 @@ please see :ref:`Structs Reference ` for more details. As with kernel parameters, all attributes of a struct must have valid type hints at class definition time. +.. _Compilation Model: Compilation Model ----------------- diff --git a/docs/faq.rst b/docs/faq.rst index 81428f640..2ab447215 100644 --- a/docs/faq.rst +++ b/docs/faq.rst @@ -35,7 +35,7 @@ Does Warp support all of the Python language? No, Warp supports a subset of Python that maps well to the GPU. Our goal is to not have any performance cliffs so that users can expect consistently good behavior from kernels that is close to native code. -Examples of unsupported concepts that don’t map well to the GPU are +Examples of unsupported concepts that don't map well to the GPU are dynamic types, list comprehensions, exceptions, garbage collection, etc. When should I call ``wp.synchronize()``? @@ -43,7 +43,7 @@ When should I call ``wp.synchronize()``? One of the common sources of confusion for new users is when calls to ``wp.synchronize()`` are necessary. The answer is “almost never”! -Synchronization is quite expensive, and should generally be avoided +Synchronization is quite expensive and should generally be avoided unless necessary. Warp naturally takes care of synchronization between operations (e.g.: kernel launches, device memory copies). @@ -83,14 +83,15 @@ and :ref:`synchronization guidance `. What happens when you differentiate a function like ``wp.abs(x)``? ------------------------------------------------------------------ -Non-smooth functions such as ``y=|x|`` do not have a single unique -gradient at ``x=0``, rather they have what is known as a -``subgradient``, which is formally the convex hull of directional +Non-smooth functions such as :math:`y=|x|` do not have a single unique +gradient at :math:`x=0`, rather they have what is known as a +*subgradient*, which is formally the convex hull of directional derivatives at that point. The way that Warp (and most auto-differentiation frameworks) handles these points is to pick an arbitrary gradient from this set, e.g.: for ``wp.abs()``, it will arbitrarily choose the gradient to be 1.0 at the origin. You can find -the implementation for these functions in ``warp/native/builtin.h``. +the implementation for these functions in +`warp/native/builtin.h `_. Most optimizers (particularly ones that exploit stochasticity), are not sensitive to the choice of which gradient to use from the subgradient, @@ -107,15 +108,15 @@ conventions of PyTorch and use aliases such as ``cuda:0``, ``cuda:1``, Should I switch to Warp over IsaacGym/PhysX? ---------------------------------------------- -Warp is not a replacement for IsaacGym, IsaacSim, or PhysX - while Warp -does offer some physical simulation capabilities this is primarily aimed +Warp is not a replacement for IsaacGym, IsaacSim, or PhysX—while Warp +does offer some physical simulation capabilities, this is primarily aimed at developers who need differentiable physics, rather than a fully featured physics engine. Warp is also integrated with IsaacGym and is great for performing auxiliary tasks such as reward and observation computations for reinforcement learning. Why aren't assignments to Warp arrays supported outside of kernels? -------------------------------------------------------------------- +------------------------------------------------------------------------ For best performance, reading and writing data that is living on the GPU can only be performed inside Warp CUDA kernels. Otherwise individual element accesses @@ -123,9 +124,9 @@ such as ``array[i] = 1.0`` in Python scope would require prohibitively slow devi synchronization and copies. We recommend to either initialize Warp arrays from other native arrays -(e.g.: Python list, NumPy array, ...) or by launching a kernel to set its values. +(Python lists, NumPy arrays, etc.) or by launching a kernel to set its values. -For the common use case of wanting to fill an array with a given value, we +For the common use case of filling an array with a given value, we also support the following forms: - ``wp.full(8, 1.23, dtype=float)``: initializes a new array of 8 float values set diff --git a/docs/limitations.rst b/docs/limitations.rst index 6e6304744..45da94128 100644 --- a/docs/limitations.rst +++ b/docs/limitations.rst @@ -40,9 +40,9 @@ This is not always possible for kernels launched with multi-dimensional grid bou `hardware limitations `_ on CUDA block dimensions. -Warp will automatically fallback to using +Warp will automatically fall back to using `grid-stride loops `_ when -it is not possible for a CUDA thread to process only one element from the Warp grid +it is not possible for a CUDA thread to process only one element from the Warp grid. When this happens, some CUDA threads may process more than one element from the Warp grid. Users can also set the ``max_blocks`` parameter to fine-tune the grid-striding behavior of kernels, even for kernels that are otherwise able to process one Warp-grid element per CUDA thread. @@ -104,7 +104,7 @@ The behavior of the modulus operator in a Warp kernel follows that of C++11: The Power Operator """""""""""""" -The power operator (``**``) in Warp kernels only works on floating-point numbers (also see :func:`wp.pow `). +The power operator (``**``) in Warp kernels only works on floating-point numbers (also see :func:`wp.pow() `). In Python, the power operator can also be used on integers. Inverse Sine and Cosine diff --git a/docs/modules/runtime.rst b/docs/modules/runtime.rst index bf1eb9724..2816edc62 100644 --- a/docs/modules/runtime.rst +++ b/docs/modules/runtime.rst @@ -15,23 +15,30 @@ Kernels are launched with the :func:`wp.launch() ` function on a specifi wp.launch(simple_kernel, dim=1024, inputs=[a, b, c], device="cuda") -Kernels may be launched with multi-dimensional grid bounds. In this case threads are not assigned a single index, +Note that all the kernel inputs must live on the target device or a runtime exception will be raised. +Kernels may be launched with multi-dimensional grid bounds. In this case, threads are not assigned a single index, but a coordinate in an n-dimensional grid, e.g.:: wp.launch(complex_kernel, dim=(128, 128, 3), ...) -Launches a 3D grid of threads with dimension 128 x 128 x 3. To retrieve the 3D index for each thread use the following syntax:: +Launches a 3D grid of threads with dimension 128 x 128 x 3. To retrieve the 3D index for each thread, use the following syntax:: i,j,k = wp.tid() .. note:: - Currently kernels launched on CPU devices will be executed in serial. + Currently, kernels launched on CPU devices will be executed in serial. Kernels launched on CUDA devices will be launched in parallel with a fixed block-size. -.. note:: - Note that all the kernel inputs must live on the target device, or a runtime exception will be raised. +In the Warp :ref:`Compilation Model`, kernels are just-in-time compiled into dynamic libraries and PTX using +C++/CUDA as an intermediate representation. +To avoid excessive runtime recompilation of kernel code, these files are stored in a cache directory +named with a module-dependent hash to allow for the reuse of previously compiled modules. +The location of the kernel cache is printed when Warp is initialized. +:func:`wp.clear_kernel_cache() ` can be used to clear the kernel cache of previously +generated compilation artifacts as Warp does not automatically try to keep the cache below a certain size. .. autofunction:: launch +.. autofunction:: clear_kernel_cache .. _Runtime Kernel Creation: diff --git a/warp/__init__.py b/warp/__init__.py index c70aae804..2822e79f9 100644 --- a/warp/__init__.py +++ b/warp/__init__.py @@ -10,7 +10,7 @@ # for autocomplete on builtins # from warp.stubs import * -from warp.types import array, array1d, array2d, array3d, array4d, constant +from warp.types import array, array1d, array2d, array3d, array4d, constant, from_ptr from warp.types import indexedarray, indexedarray1d, indexedarray2d, indexedarray3d, indexedarray4d from warp.fabric import fabricarray, fabricarrayarray, indexedfabricarray, indexedfabricarrayarray @@ -31,9 +31,9 @@ from warp.types import bvh_query_t, hash_grid_query_t, mesh_query_aabb_t, mesh_query_point_t, mesh_query_ray_t # device-wide gemms -from warp.types import matmul, adj_matmul, batched_matmul, adj_batched_matmul, from_ptr +from warp.types import matmul, adj_matmul, batched_matmul, adj_batched_matmul -# deprecated +# discouraged, users should use wp.types.vector, wp.types.matrix from warp.types import vector as vec from warp.types import matrix as mat @@ -99,6 +99,8 @@ from warp.dlpack import from_dlpack, to_dlpack +from warp.build import clear_kernel_cache + from warp.constants import * from . import builtins diff --git a/warp/build.py b/warp/build.py index eb3bc2244..7eee8e291 100644 --- a/warp/build.py +++ b/warp/build.py @@ -71,8 +71,12 @@ def init_kernel_cache(path=None): os.makedirs(warp.config.kernel_cache_dir, exist_ok=True) -def clear_kernel_cache(): - """Clear the kernel cache.""" +def clear_kernel_cache() -> None: + """Clear the kernel cache directory of previously generated source code and compiler artifacts. + + Only directories beginning with ``wp_`` will be deleted. + This function only clears the cache for the current Warp version. + """ warp.context.init() diff --git a/warp/examples/benchmarks/benchmark_cloth_warp.py b/warp/examples/benchmarks/benchmark_cloth_warp.py index 9b52a42de..c55dc4a3a 100644 --- a/warp/examples/benchmarks/benchmark_cloth_warp.py +++ b/warp/examples/benchmarks/benchmark_cloth_warp.py @@ -7,7 +7,7 @@ import warp as wp -wp.build.clear_kernel_cache() +wp.clear_kernel_cache() @wp.kernel diff --git a/warp/examples/benchmarks/benchmark_launches.py b/warp/examples/benchmarks/benchmark_launches.py index a21ab9e74..daa0b769e 100644 --- a/warp/examples/benchmarks/benchmark_launches.py +++ b/warp/examples/benchmarks/benchmark_launches.py @@ -128,7 +128,7 @@ def ksz(s: Sz): tid = wp.tid() # noqa: F841 -wp.build.clear_kernel_cache() +wp.clear_kernel_cache() devices = wp.get_devices() num_launches = 100000 diff --git a/warp/stubs.py b/warp/stubs.py index 1218fe908..bbc2d1134 100644 --- a/warp/stubs.py +++ b/warp/stubs.py @@ -23,7 +23,7 @@ IndexedFabricArray = Generic[DType] -from warp.types import array, array1d, array2d, array3d, array4d, constant +from warp.types import array, array1d, array2d, array3d, array4d, constant, from_ptr from warp.types import indexedarray, indexedarray1d, indexedarray2d, indexedarray3d, indexedarray4d from warp.fabric import fabricarray, fabricarrayarray, indexedfabricarray, indexedfabricarrayarray @@ -42,7 +42,7 @@ from warp.types import Bvh, Mesh, HashGrid, Volume, MarchingCubes from warp.types import bvh_query_t, hash_grid_query_t, mesh_query_aabb_t, mesh_query_point_t, mesh_query_ray_t -from warp.types import matmul, adj_matmul, batched_matmul, adj_batched_matmul, from_ptr +from warp.types import matmul, adj_matmul, batched_matmul, adj_batched_matmul from warp.types import vector as vec from warp.types import matrix as mat @@ -108,6 +108,8 @@ from warp.dlpack import from_dlpack, to_dlpack +from warp.build import clear_kernel_cache + from warp.constants import * from . import builtins diff --git a/warp/tests/disabled_kinematics.py b/warp/tests/disabled_kinematics.py index 7451e55af..c9caf3399 100644 --- a/warp/tests/disabled_kinematics.py +++ b/warp/tests/disabled_kinematics.py @@ -233,5 +233,5 @@ class TestKinematics(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2, failfast=False) diff --git a/warp/tests/test_adam.py b/warp/tests/test_adam.py index 1fa14ca81..42199df9a 100644 --- a/warp/tests/test_adam.py +++ b/warp/tests/test_adam.py @@ -151,5 +151,5 @@ class TestAdam(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_arithmetic.py b/warp/tests/test_arithmetic.py index 57359d132..4950adfa5 100644 --- a/warp/tests/test_arithmetic.py +++ b/warp/tests/test_arithmetic.py @@ -1084,5 +1084,5 @@ class TestArithmetic(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2, failfast=False) diff --git a/warp/tests/test_array.py b/warp/tests/test_array.py index 049753c9c..c40254434 100644 --- a/warp/tests/test_array.py +++ b/warp/tests/test_array.py @@ -2411,5 +2411,5 @@ def test_array_new_del(self): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_array_reduce.py b/warp/tests/test_array_reduce.py index cd331ed33..adf098ad9 100644 --- a/warp/tests/test_array_reduce.py +++ b/warp/tests/test_array_reduce.py @@ -144,5 +144,5 @@ class TestArrayReduce(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_async.py b/warp/tests/test_async.py index f5dec8c08..f5d541f51 100644 --- a/warp/tests/test_async.py +++ b/warp/tests/test_async.py @@ -662,5 +662,5 @@ def test_func( # value_offset=0)) if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_atomic.py b/warp/tests/test_atomic.py index 4cfa68fac..2c8142b38 100644 --- a/warp/tests/test_atomic.py +++ b/warp/tests/test_atomic.py @@ -135,5 +135,5 @@ class TestAtomic(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_bool.py b/warp/tests/test_bool.py index 7dd031c55..aff1c9eec 100644 --- a/warp/tests/test_bool.py +++ b/warp/tests/test_bool.py @@ -208,5 +208,5 @@ class TestBool(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_builtins_resolution.py b/warp/tests/test_builtins_resolution.py index fcca3b6be..8a72d59e7 100644 --- a/warp/tests/test_builtins_resolution.py +++ b/warp/tests/test_builtins_resolution.py @@ -1286,5 +1286,5 @@ def test_vec4_float_args_precision(self): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_bvh.py b/warp/tests/test_bvh.py index 2473529f6..9971b708e 100644 --- a/warp/tests/test_bvh.py +++ b/warp/tests/test_bvh.py @@ -163,5 +163,5 @@ def test_bvh_new_del(self): add_function_test(TestBvh, "test_bvh_ray", test_bvh_query_ray, devices=devices) if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_closest_point_edge_edge.py b/warp/tests/test_closest_point_edge_edge.py index 49bc5935b..37d5c8acd 100644 --- a/warp/tests/test_closest_point_edge_edge.py +++ b/warp/tests/test_closest_point_edge_edge.py @@ -223,5 +223,5 @@ class TestClosestPointEdgeEdgeMethods(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_codegen.py b/warp/tests/test_codegen.py index 8d391a9b5..bbe09ce6f 100644 --- a/warp/tests/test_codegen.py +++ b/warp/tests/test_codegen.py @@ -558,5 +558,5 @@ class TestCodeGen(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2, failfast=True) diff --git a/warp/tests/test_compile_consts.py b/warp/tests/test_compile_consts.py index b3f3b6411..ddc7734d4 100644 --- a/warp/tests/test_compile_consts.py +++ b/warp/tests/test_compile_consts.py @@ -213,5 +213,5 @@ def test_constant_math(self): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_conditional.py b/warp/tests/test_conditional.py index 06b00388c..2626421da 100644 --- a/warp/tests/test_conditional.py +++ b/warp/tests/test_conditional.py @@ -240,5 +240,5 @@ class TestConditional(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_copy.py b/warp/tests/test_copy.py index acd61342b..c0f6fd1d8 100644 --- a/warp/tests/test_copy.py +++ b/warp/tests/test_copy.py @@ -226,5 +226,5 @@ class TestCopy(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_ctypes.py b/warp/tests/test_ctypes.py index 37d8555c6..0e992b10a 100644 --- a/warp/tests/test_ctypes.py +++ b/warp/tests/test_ctypes.py @@ -626,5 +626,5 @@ class TestCTypes(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_dense.py b/warp/tests/test_dense.py index df77fbbf2..fde47ac52 100644 --- a/warp/tests/test_dense.py +++ b/warp/tests/test_dense.py @@ -61,5 +61,5 @@ class TestDense(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_devices.py b/warp/tests/test_devices.py index 8e821912e..df8010293 100644 --- a/warp/tests/test_devices.py +++ b/warp/tests/test_devices.py @@ -85,5 +85,5 @@ def test_devices_unmap_imaginary_device(self): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_dlpack.py b/warp/tests/test_dlpack.py index fafe53a01..ad45806ab 100644 --- a/warp/tests/test_dlpack.py +++ b/warp/tests/test_dlpack.py @@ -524,5 +524,5 @@ class TestDLPack(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_examples.py b/warp/tests/test_examples.py index 2950f6822..5a5dc474e 100644 --- a/warp/tests/test_examples.py +++ b/warp/tests/test_examples.py @@ -398,6 +398,6 @@ class TestFemDiffusionExamples(unittest.TestCase): if __name__ == "__main__": # force rebuild of all kernels - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2, failfast=True) diff --git a/warp/tests/test_fabricarray.py b/warp/tests/test_fabricarray.py index b691ae6ba..0bf0f41ff 100644 --- a/warp/tests/test_fabricarray.py +++ b/warp/tests/test_fabricarray.py @@ -951,5 +951,5 @@ def test_fabricarray_new_del(self): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_fast_math.py b/warp/tests/test_fast_math.py index 974ac86db..35b4692b0 100644 --- a/warp/tests/test_fast_math.py +++ b/warp/tests/test_fast_math.py @@ -56,5 +56,5 @@ def test_fast_math_cpu(self): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_fem.py b/warp/tests/test_fem.py index aa48a8626..2cc1f8bc0 100644 --- a/warp/tests/test_fem.py +++ b/warp/tests/test_fem.py @@ -1294,5 +1294,5 @@ class TestFemShapeFunctions(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_fp16.py b/warp/tests/test_fp16.py index da43b552e..506fc40cb 100644 --- a/warp/tests/test_fp16.py +++ b/warp/tests/test_fp16.py @@ -124,5 +124,5 @@ class TestFp16(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_func.py b/warp/tests/test_func.py index 06229fe6a..1a771be0a 100644 --- a/warp/tests/test_func.py +++ b/warp/tests/test_func.py @@ -332,5 +332,5 @@ def test_native_function_error_resolution(self): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_generics.py b/warp/tests/test_generics.py index a01ab31ad..ed769338c 100644 --- a/warp/tests/test_generics.py +++ b/warp/tests/test_generics.py @@ -592,5 +592,5 @@ class TestGenerics(unittest.TestCase): add_function_test(TestGenerics, "test_type_attribute_error", test_type_attribute_error, devices=devices) if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_grad.py b/warp/tests/test_grad.py index 1a4958037..f4b1b134b 100644 --- a/warp/tests/test_grad.py +++ b/warp/tests/test_grad.py @@ -881,5 +881,5 @@ class TestGrad(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2, failfast=False) diff --git a/warp/tests/test_grad_customs.py b/warp/tests/test_grad_customs.py index cd65a02de..5a08aad33 100644 --- a/warp/tests/test_grad_customs.py +++ b/warp/tests/test_grad_customs.py @@ -327,5 +327,5 @@ def test_wrapped_docstring(self): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2, failfast=False) diff --git a/warp/tests/test_hash_grid.py b/warp/tests/test_hash_grid.py index 4f5a236b0..c6133f6aa 100644 --- a/warp/tests/test_hash_grid.py +++ b/warp/tests/test_hash_grid.py @@ -209,5 +209,5 @@ def test_hashgrid_new_del(self): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2, failfast=False) diff --git a/warp/tests/test_import.py b/warp/tests/test_import.py index 1eb7e166a..ab08580fa 100644 --- a/warp/tests/test_import.py +++ b/warp/tests/test_import.py @@ -33,5 +33,5 @@ class TestImport(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_indexedarray.py b/warp/tests/test_indexedarray.py index aaf98e58b..1c57d6bb8 100644 --- a/warp/tests/test_indexedarray.py +++ b/warp/tests/test_indexedarray.py @@ -1128,5 +1128,5 @@ class TestIndexedArray(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_intersect.py b/warp/tests/test_intersect.py index ea5e55309..7e4dd9bf1 100644 --- a/warp/tests/test_intersect.py +++ b/warp/tests/test_intersect.py @@ -61,5 +61,5 @@ class TestIntersect(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2, failfast=False) diff --git a/warp/tests/test_jax.py b/warp/tests/test_jax.py index 6aeec700c..7cb24db3b 100644 --- a/warp/tests/test_jax.py +++ b/warp/tests/test_jax.py @@ -301,5 +301,5 @@ class TestJax(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_large.py b/warp/tests/test_large.py index d8189a5c4..a926701bf 100644 --- a/warp/tests/test_large.py +++ b/warp/tests/test_large.py @@ -165,5 +165,5 @@ class TestLarge(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_launch.py b/warp/tests/test_launch.py index 23af4eed4..85c1d22eb 100644 --- a/warp/tests/test_launch.py +++ b/warp/tests/test_launch.py @@ -348,5 +348,5 @@ class TestLaunch(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_lerp.py b/warp/tests/test_lerp.py index e368ba8da..7fe48bbaf 100644 --- a/warp/tests/test_lerp.py +++ b/warp/tests/test_lerp.py @@ -213,5 +213,5 @@ class TestLerp(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_linear_solvers.py b/warp/tests/test_linear_solvers.py index 3e28a0868..ac3103040 100644 --- a/warp/tests/test_linear_solvers.py +++ b/warp/tests/test_linear_solvers.py @@ -185,5 +185,5 @@ class TestLinearSolvers(unittest.TestCase): add_function_test(TestLinearSolvers, "test_gmres", test_gmres, devices=devices) if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_lvalue.py b/warp/tests/test_lvalue.py index d62f3b967..e4c07753a 100644 --- a/warp/tests/test_lvalue.py +++ b/warp/tests/test_lvalue.py @@ -415,5 +415,5 @@ def test_swizzle_error_invalid_attribute(self): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_marching_cubes.py b/warp/tests/test_marching_cubes.py index a19791db2..907e1c329 100644 --- a/warp/tests/test_marching_cubes.py +++ b/warp/tests/test_marching_cubes.py @@ -62,5 +62,5 @@ def test_marching_cubes_new_del(self): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_mat.py b/warp/tests/test_mat.py index 15de2181e..43623f151 100644 --- a/warp/tests/test_mat.py +++ b/warp/tests/test_mat.py @@ -1795,5 +1795,5 @@ def test_tpl_ops_with_anon(self): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2, failfast=True) diff --git a/warp/tests/test_mat_lite.py b/warp/tests/test_mat_lite.py index 8c61f617d..990785206 100644 --- a/warp/tests/test_mat_lite.py +++ b/warp/tests/test_mat_lite.py @@ -109,5 +109,5 @@ class TestMatLite(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2, failfast=True) diff --git a/warp/tests/test_mat_scalar_ops.py b/warp/tests/test_mat_scalar_ops.py index 9d2eb86cc..67b6c0c71 100644 --- a/warp/tests/test_mat_scalar_ops.py +++ b/warp/tests/test_mat_scalar_ops.py @@ -2901,5 +2901,5 @@ class TestMatScalarOps(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2, failfast=True) diff --git a/warp/tests/test_math.py b/warp/tests/test_math.py index bd76851db..62425dacd 100644 --- a/warp/tests/test_math.py +++ b/warp/tests/test_math.py @@ -120,5 +120,5 @@ def test_mat_type(self): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_matmul.py b/warp/tests/test_matmul.py index c25c7266e..7f788618a 100644 --- a/warp/tests/test_matmul.py +++ b/warp/tests/test_matmul.py @@ -494,5 +494,5 @@ class TestMatmul(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2, failfast=False) diff --git a/warp/tests/test_matmul_lite.py b/warp/tests/test_matmul_lite.py index a3c1a9a0f..ed50d4039 100644 --- a/warp/tests/test_matmul_lite.py +++ b/warp/tests/test_matmul_lite.py @@ -404,5 +404,5 @@ class TestMatmulLite(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2, failfast=False) diff --git a/warp/tests/test_mempool.py b/warp/tests/test_mempool.py index 8f4002b64..23751bc36 100644 --- a/warp/tests/test_mempool.py +++ b/warp/tests/test_mempool.py @@ -182,5 +182,5 @@ class TestMempool(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_mesh.py b/warp/tests/test_mesh.py index a5f62e11b..94675c75c 100644 --- a/warp/tests/test_mesh.py +++ b/warp/tests/test_mesh.py @@ -280,5 +280,5 @@ def test_mesh_new_del(self): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_mesh_query_aabb.py b/warp/tests/test_mesh_query_aabb.py index a80ea38d1..53c666003 100644 --- a/warp/tests/test_mesh_query_aabb.py +++ b/warp/tests/test_mesh_query_aabb.py @@ -222,5 +222,5 @@ def kernel_fn( ) if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_mesh_query_point.py b/warp/tests/test_mesh_query_point.py index 9d6b7ff00..47a97b057 100644 --- a/warp/tests/test_mesh_query_point.py +++ b/warp/tests/test_mesh_query_point.py @@ -686,5 +686,5 @@ def kernel_fn( add_function_test(TestMeshQueryPoint, "test_adj_mesh_query_point", test_adj_mesh_query_point, devices=devices) if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_mesh_query_ray.py b/warp/tests/test_mesh_query_ray.py index 56be12803..ed7dd20ad 100644 --- a/warp/tests/test_mesh_query_ray.py +++ b/warp/tests/test_mesh_query_ray.py @@ -286,5 +286,5 @@ def kernel_fn( if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_mlp.py b/warp/tests/test_mlp.py index 51c5b6d5b..e8397126c 100644 --- a/warp/tests/test_mlp.py +++ b/warp/tests/test_mlp.py @@ -270,5 +270,5 @@ class TestMLP(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2, failfast=False) diff --git a/warp/tests/test_model.py b/warp/tests/test_model.py index 72c1f54ed..cf3f547b9 100644 --- a/warp/tests/test_model.py +++ b/warp/tests/test_model.py @@ -104,5 +104,5 @@ def test_add_edges(self): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_module_hashing.py b/warp/tests/test_module_hashing.py index 281d53ae0..cca964326 100644 --- a/warp/tests/test_module_hashing.py +++ b/warp/tests/test_module_hashing.py @@ -107,5 +107,5 @@ class TestModuleHashing(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_modules_lite.py b/warp/tests/test_modules_lite.py index fbe941684..5f43e08af 100644 --- a/warp/tests/test_modules_lite.py +++ b/warp/tests/test_modules_lite.py @@ -32,5 +32,5 @@ def test_module_lite_options(self): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_multigpu.py b/warp/tests/test_multigpu.py index a1df19e55..55459fbaf 100644 --- a/warp/tests/test_multigpu.py +++ b/warp/tests/test_multigpu.py @@ -157,5 +157,5 @@ def test_multigpu_pingpong_streams(self): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2, failfast=False) diff --git a/warp/tests/test_noise.py b/warp/tests/test_noise.py index 2bb02dadb..1f4ffd171 100644 --- a/warp/tests/test_noise.py +++ b/warp/tests/test_noise.py @@ -240,5 +240,5 @@ class TestNoise(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_operators.py b/warp/tests/test_operators.py index d45e5bae1..5356029f9 100644 --- a/warp/tests/test_operators.py +++ b/warp/tests/test_operators.py @@ -244,5 +244,5 @@ class TestOperators(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_options.py b/warp/tests/test_options.py index 12e2471e1..692bc56eb 100644 --- a/warp/tests/test_options.py +++ b/warp/tests/test_options.py @@ -117,5 +117,5 @@ class TestOptions(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_peer.py b/warp/tests/test_peer.py index 3959949fa..0674458a7 100644 --- a/warp/tests/test_peer.py +++ b/warp/tests/test_peer.py @@ -127,5 +127,5 @@ class TestPeer(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_pinned.py b/warp/tests/test_pinned.py index 53e3370c4..f40c9faf6 100644 --- a/warp/tests/test_pinned.py +++ b/warp/tests/test_pinned.py @@ -72,5 +72,5 @@ class TestPinned(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_print.py b/warp/tests/test_print.py index 138b3d62a..871f7de5b 100644 --- a/warp/tests/test_print.py +++ b/warp/tests/test_print.py @@ -48,5 +48,5 @@ class TestPrint(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_quat.py b/warp/tests/test_quat.py index 00216cfb7..899722e6d 100644 --- a/warp/tests/test_quat.py +++ b/warp/tests/test_quat.py @@ -2080,5 +2080,5 @@ class TestQuat(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_rand.py b/warp/tests/test_rand.py index 948ce1c3b..30823bb95 100644 --- a/warp/tests/test_rand.py +++ b/warp/tests/test_rand.py @@ -320,5 +320,5 @@ class TestRand(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_reload.py b/warp/tests/test_reload.py index e4d0e8c8f..626955dd2 100644 --- a/warp/tests/test_reload.py +++ b/warp/tests/test_reload.py @@ -203,5 +203,5 @@ class TestReload(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2, failfast=False) diff --git a/warp/tests/test_rounding.py b/warp/tests/test_rounding.py index 17a95ba60..a434b5d06 100644 --- a/warp/tests/test_rounding.py +++ b/warp/tests/test_rounding.py @@ -173,5 +173,5 @@ class TestRounding(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_runlength_encode.py b/warp/tests/test_runlength_encode.py index 1525a6973..347d2451d 100644 --- a/warp/tests/test_runlength_encode.py +++ b/warp/tests/test_runlength_encode.py @@ -184,5 +184,5 @@ def test_runlength_encode_error_run_count_device_mismatch(self): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_sim_grad.py b/warp/tests/test_sim_grad.py index 42ef3ef2d..7b6d677e7 100644 --- a/warp/tests/test_sim_grad.py +++ b/warp/tests/test_sim_grad.py @@ -237,5 +237,5 @@ def test_fn(self, device, jt_type=jt_type, int_type=int_type): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2, failfast=True) diff --git a/warp/tests/test_sim_kinematics.py b/warp/tests/test_sim_kinematics.py index 7112a38d8..2a9a5aae9 100644 --- a/warp/tests/test_sim_kinematics.py +++ b/warp/tests/test_sim_kinematics.py @@ -85,5 +85,5 @@ class TestSimKinematics(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2, failfast=True) diff --git a/warp/tests/test_smoothstep.py b/warp/tests/test_smoothstep.py index 1fa98f68e..73bb393c6 100644 --- a/warp/tests/test_smoothstep.py +++ b/warp/tests/test_smoothstep.py @@ -162,5 +162,5 @@ class TestSmoothstep(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_sparse.py b/warp/tests/test_sparse.py index 513a9719e..caf4e423b 100644 --- a/warp/tests/test_sparse.py +++ b/warp/tests/test_sparse.py @@ -529,5 +529,5 @@ def test_bsr_copy_scale(self): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_spatial.py b/warp/tests/test_spatial.py index 46e69febe..4eb21a8c8 100644 --- a/warp/tests/test_spatial.py +++ b/warp/tests/test_spatial.py @@ -2142,5 +2142,5 @@ class TestSpatial(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_special_values.py b/warp/tests/test_special_values.py index 28f37d496..2917843c8 100644 --- a/warp/tests/test_special_values.py +++ b/warp/tests/test_special_values.py @@ -358,5 +358,5 @@ class TestSpecialValues(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2, failfast=False) diff --git a/warp/tests/test_streams.py b/warp/tests/test_streams.py index 12a96f0f6..67c7f0560 100644 --- a/warp/tests/test_streams.py +++ b/warp/tests/test_streams.py @@ -490,5 +490,5 @@ def test_event_new_del(self): add_function_test(TestStreams, "test_event_elapsed_time", test_event_elapsed_time, devices=devices) if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_struct.py b/warp/tests/test_struct.py index db15f8419..81594be27 100644 --- a/warp/tests/test_struct.py +++ b/warp/tests/test_struct.py @@ -704,5 +704,5 @@ def test_nested_vec_assignment(self): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_tape.py b/warp/tests/test_tape.py index 0f5e54c48..596649234 100644 --- a/warp/tests/test_tape.py +++ b/warp/tests/test_tape.py @@ -167,5 +167,5 @@ def test_tape_no_nested_tapes(self): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_torch.py b/warp/tests/test_torch.py index f6588639a..4b3732ab2 100644 --- a/warp/tests/test_torch.py +++ b/warp/tests/test_torch.py @@ -737,5 +737,5 @@ class TestTorch(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_transient_module.py b/warp/tests/test_transient_module.py index 9db31edfe..88a81918a 100644 --- a/warp/tests/test_transient_module.py +++ b/warp/tests/test_transient_module.py @@ -81,5 +81,5 @@ class TestTransientModule(unittest.TestCase): add_function_test(TestTransientModule, "test_transient_module", test_transient_module, devices=devices) if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_types.py b/warp/tests/test_types.py index abe1ac290..51f5f99b7 100644 --- a/warp/tests/test_types.py +++ b/warp/tests/test_types.py @@ -550,5 +550,5 @@ def test_conversions(warp_type, np_type): add_function_test(TestTypes, f"test_vector_{dtype.__name__}", test_vector, devices=devices, dtype=dtype) if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_utils.py b/warp/tests/test_utils.py index 763153a95..74ce0bbe0 100644 --- a/warp/tests/test_utils.py +++ b/warp/tests/test_utils.py @@ -484,5 +484,5 @@ def test_scoped_timer(self): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_vec.py b/warp/tests/test_vec.py index 86859ff00..21fbbc03e 100644 --- a/warp/tests/test_vec.py +++ b/warp/tests/test_vec.py @@ -1258,5 +1258,5 @@ def test_tpl_ops_with_anon(self): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2, failfast=True) diff --git a/warp/tests/test_vec_lite.py b/warp/tests/test_vec_lite.py index ba4064103..44ac0b1b3 100644 --- a/warp/tests/test_vec_lite.py +++ b/warp/tests/test_vec_lite.py @@ -67,5 +67,5 @@ class TestVecLite(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2, failfast=True) diff --git a/warp/tests/test_vec_scalar_ops.py b/warp/tests/test_vec_scalar_ops.py index d1dc4b4ed..5ce1b7861 100644 --- a/warp/tests/test_vec_scalar_ops.py +++ b/warp/tests/test_vec_scalar_ops.py @@ -2112,5 +2112,5 @@ class TestVecScalarOps(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2, failfast=True) diff --git a/warp/tests/test_verify_fp.py b/warp/tests/test_verify_fp.py index 993f0d2d7..21ab08e43 100644 --- a/warp/tests/test_verify_fp.py +++ b/warp/tests/test_verify_fp.py @@ -88,5 +88,5 @@ class TestVerifyFP(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_volume.py b/warp/tests/test_volume.py index 727d24071..9901d5ff9 100644 --- a/warp/tests/test_volume.py +++ b/warp/tests/test_volume.py @@ -960,5 +960,5 @@ def test_volume_new_del(self): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/test_volume_write.py b/warp/tests/test_volume_write.py index 9d1a9f1bc..8ad026143 100644 --- a/warp/tests/test_volume_write.py +++ b/warp/tests/test_volume_write.py @@ -334,5 +334,5 @@ class TestVolumeWrite(unittest.TestCase): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() unittest.main(verbosity=2) diff --git a/warp/tests/unittest_serial.py b/warp/tests/unittest_serial.py index 8cf822191..1323d0f16 100644 --- a/warp/tests/unittest_serial.py +++ b/warp/tests/unittest_serial.py @@ -15,7 +15,7 @@ def run_suite() -> bool: """Run a test suite""" # force rebuild of all kernels - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() print("Cleared Warp kernel cache") runner = unittest.TextTestRunner(verbosity=2, failfast=True) diff --git a/warp/tests/walkthrough_debug.py b/warp/tests/walkthrough_debug.py index 9b590d86b..2c26d963c 100644 --- a/warp/tests/walkthrough_debug.py +++ b/warp/tests/walkthrough_debug.py @@ -81,5 +81,5 @@ def example_breakpoint(n: int): if __name__ == "__main__": - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() wp.launch(example_breakpoint, dim=1, inputs=[10], device="cpu") diff --git a/warp/thirdparty/unittest_parallel.py b/warp/thirdparty/unittest_parallel.py index 286e27fb5..ac117c48f 100644 --- a/warp/thirdparty/unittest_parallel.py +++ b/warp/thirdparty/unittest_parallel.py @@ -173,7 +173,7 @@ def main(argv=None): import warp as wp # NVIDIA Modification # Clear the Warp cache (NVIDIA Modification) - wp.build.clear_kernel_cache() + wp.clear_kernel_cache() print("Cleared Warp kernel cache") # Create the temporary directory (for coverage files)