From c3e2195466a2ce00d3ccc548686c3109e905a183 Mon Sep 17 00:00:00 2001 From: Christopher Crouzet Date: Wed, 12 Jun 2024 09:14:18 +1200 Subject: [PATCH 1/5] Fix array constructors not initializing Warp --- warp/types.py | 28 ++++++++++++++++------------ 1 file changed, 16 insertions(+), 12 deletions(-) diff --git a/warp/types.py b/warp/types.py index 7946c2427..d277d96c0 100644 --- a/warp/types.py +++ b/warp/types.py @@ -1650,9 +1650,10 @@ def _init_from_data(self, data, dtype, shape, device, copy, pinned): try: # Performance note: try first, ask questions later device = warp.context.runtime.get_device(device) - except: - warp.context.init() - raise + except Exception: + # Fallback to using the public API for retrieving the device, + # which takes take of initializing Warp if needed. + device = warp.context.get_device(device) if device.is_cuda: desc = data.__cuda_array_interface__ @@ -1779,9 +1780,10 @@ def _init_from_data(self, data, dtype, shape, device, copy, pinned): try: # Performance note: try first, ask questions later device = warp.context.runtime.get_device(device) - except: - warp.context.init() - raise + except Exception: + # Fallback to using the public API for retrieving the device, + # which takes take of initializing Warp if needed. + device = warp.context.get_device(device) if device.is_cpu and not copy and not pinned: # reference numpy memory directly @@ -1805,9 +1807,10 @@ def _init_from_ptr(self, ptr, dtype, shape, strides, capacity, device, pinned, d try: # Performance note: try first, ask questions later device = warp.context.runtime.get_device(device) - except: - warp.context.init() - raise + except Exception: + # Fallback to using the public API for retrieving the device, + # which takes take of initializing Warp if needed. + device = warp.context.get_device(device) check_array_shape(shape) ndim = len(shape) @@ -1852,9 +1855,10 @@ def _init_new(self, dtype, shape, strides, device, pinned): try: # Performance note: try first, ask questions later device = warp.context.runtime.get_device(device) - except: - warp.context.init() - raise + except Exception: + # Fallback to using the public API for retrieving the device, + # which takes take of initializing Warp if needed. + device = warp.context.get_device(device) check_array_shape(shape) ndim = len(shape) From ef4456ea8b7d7c130b4b45250d203e84837d11c7 Mon Sep 17 00:00:00 2001 From: Christopher Crouzet Date: Wed, 12 Jun 2024 11:06:08 +1200 Subject: [PATCH 2/5] Fix omission of device retrieval --- warp/context.py | 3 +++ 1 file changed, 3 insertions(+) diff --git a/warp/context.py b/warp/context.py index 9b160206f..60016e40d 100644 --- a/warp/context.py +++ b/warp/context.py @@ -3632,6 +3632,9 @@ def is_mempool_access_supported(target_device: Devicelike, peer_device: Deviceli init() + target_device = runtime.get_device(target_device) + peer_device = runtime.get_device(peer_device) + return target_device.is_mempool_supported and is_peer_access_supported(target_device, peer_device) From 86620ef700519b845186726ef1cc5d98efb89c9c Mon Sep 17 00:00:00 2001 From: Christopher Crouzet Date: Wed, 12 Jun 2024 08:04:33 +1200 Subject: [PATCH 3/5] Add tests to detect uninitialized runtime --- warp/context.py | 2 +- warp/tests/test_implicit_init.py | 354 +++++++++++++++++++++++++++++++ 2 files changed, 355 insertions(+), 1 deletion(-) create mode 100644 warp/tests/test_implicit_init.py diff --git a/warp/context.py b/warp/context.py index 60016e40d..c758e8ff2 100644 --- a/warp/context.py +++ b/warp/context.py @@ -347,7 +347,7 @@ def __repr__(self): def call_builtin(func: Function, *params) -> Tuple[bool, Any]: uses_non_warp_array_type = False - warp.context.init() + init() # Retrieve the built-in function from Warp's dll. c_func = getattr(warp.context.runtime.core, func.mangled_name) diff --git a/warp/tests/test_implicit_init.py b/warp/tests/test_implicit_init.py new file mode 100644 index 000000000..e9daef584 --- /dev/null +++ b/warp/tests/test_implicit_init.py @@ -0,0 +1,354 @@ +# Copyright (c) 2024 NVIDIA CORPORATION. All rights reserved. +# NVIDIA CORPORATION and its licensors retain all intellectual property +# and proprietary rights in and to this software, related documentation +# and any modifications thereto. Any use, reproduction, disclosure or +# distribution of this software and related documentation without an express +# license agreement from NVIDIA CORPORATION is strictly prohibited. + +import unittest + +import warp as wp +from warp.tests.unittest_utils import * + +# Array Initialization +# ------------------------------------------------------------------------------ + + +def test_array_from_data(test, device): + wp.array((1.0, 2.0, 3.0), dtype=float) + + +class TestImplicitInitArrayFromData(unittest.TestCase): + pass + + +add_function_test( + TestImplicitInitArrayFromData, + "test_array_from_data", + test_array_from_data, + check_output=False, +) + + +def test_array_from_ptr(test, device): + wp.array(ptr=0, shape=(123,), dtype=float) + + +class TestImplicitInitArrayFromPtr(unittest.TestCase): + pass + + +add_function_test( + TestImplicitInitArrayFromPtr, + "test_array_from_ptr", + test_array_from_ptr, + check_output=False, +) + + +# Builtin Call +# ------------------------------------------------------------------------------ + + +def test_builtin_call(test, device): + wp.sin(1.23) + + +class TestImplicitInitBuiltinCall(unittest.TestCase): + pass + + +add_function_test( + TestImplicitInitBuiltinCall, + "test_builtin_call", + test_builtin_call, + check_output=False, +) + + +# Devices +# ------------------------------------------------------------------------------ + + +def test_get_cuda_device_count(test, device): + wp.get_cuda_device_count() + + +class TestImplicitInitGetCudaDeviceCount(unittest.TestCase): + pass + + +add_function_test( + TestImplicitInitGetCudaDeviceCount, + "test_get_cuda_device_count", + test_get_cuda_device_count, + check_output=False, +) + + +def test_get_cuda_devices(test, device): + wp.get_cuda_devices() + + +class TestImplicitInitGetCudaDevices(unittest.TestCase): + pass + + +add_function_test( + TestImplicitInitGetCudaDevices, + "test_get_cuda_devices", + test_get_cuda_devices, + check_output=False, +) + + +def test_get_device(test, device): + wp.get_device("cpu") + + +class TestImplicitInitGetDevice(unittest.TestCase): + pass + + +add_function_test( + TestImplicitInitGetDevice, + "test_get_device", + test_get_device, + check_output=False, +) + + +def test_get_devices(test, device): + wp.get_devices() + + +class TestImplicitInitGetDevices(unittest.TestCase): + pass + + +add_function_test( + TestImplicitInitGetDevices, + "test_get_devices", + test_get_devices, + check_output=False, +) + + +def test_get_preferred_device(test, device): + wp.get_preferred_device() + + +class TestImplicitInitGetPreferredDevice(unittest.TestCase): + pass + + +add_function_test( + TestImplicitInitGetPreferredDevice, + "test_get_preferred_device", + test_get_preferred_device, + check_output=False, +) + + +def test_is_cpu_available(test, device): + wp.is_cpu_available() + + +class TestImplicitInitIsCpuAvailable(unittest.TestCase): + pass + + +add_function_test( + TestImplicitInitIsCpuAvailable, + "test_is_cpu_available", + test_is_cpu_available, + check_output=False, +) + + +def test_is_cuda_available(test, device): + wp.is_cuda_available() + + +class TestImplicitInitIsCudaAvailable(unittest.TestCase): + pass + + +add_function_test( + TestImplicitInitIsCudaAvailable, + "test_is_cuda_available", + test_is_cuda_available, + check_output=False, +) + + +def test_is_device_available(test, device): + wp.is_device_available("cpu") + + +class TestImplicitInitIsDeviceAvailable(unittest.TestCase): + pass + + +add_function_test( + TestImplicitInitIsDeviceAvailable, + "test_is_device_available", + test_is_device_available, + check_output=False, +) + + +def test_set_device(test, device): + wp.set_device("cpu") + + +class TestImplicitInitSetDevice(unittest.TestCase): + pass + + +add_function_test( + TestImplicitInitSetDevice, + "test_set_device", + test_set_device, + check_output=False, +) + + +# Launch +# ------------------------------------------------------------------------------ + + +@wp.kernel +def launch_kernel(): + pass + + +def test_launch(test, device): + wp.launch(launch_kernel, dim=1) + + +class TestImplicitInitLaunch(unittest.TestCase): + pass + + +add_function_test( + TestImplicitInitLaunch, + "test_launch", + test_launch, + check_output=False, +) + + +# Mempool +# ------------------------------------------------------------------------------ + + +def test_is_mempool_enabled(test, device): + wp.is_mempool_enabled("cpu") + + +class TestImplicitInitIsMempoolEnabled(unittest.TestCase): + pass + + +add_function_test( + TestImplicitInitIsMempoolEnabled, + "test_is_mempool_enabled", + test_is_mempool_enabled, + check_output=False, +) + + +def test_is_mempool_supported(test, device): + wp.is_mempool_supported("cpu") + + +class TestImplicitInitIsMempoolSupported(unittest.TestCase): + pass + + +add_function_test( + TestImplicitInitIsMempoolSupported, + "test_is_mempool_supported", + test_is_mempool_supported, + check_output=False, +) + + +# Mempool Access +# ------------------------------------------------------------------------------ + + +def test_is_mempool_access_enabled(test, device): + wp.is_mempool_access_enabled("cpu", "cpu") + + +class TestImplicitInitIsMempoolAccessEnabled(unittest.TestCase): + pass + + +add_function_test( + TestImplicitInitIsMempoolAccessEnabled, + "test_is_mempool_access_enabled", + test_is_mempool_access_enabled, + check_output=False, +) + + +def test_is_mempool_access_supported(test, device): + wp.is_mempool_access_supported("cpu", "cpu") + + +class TestImplicitInitIsMempoolAccessSupported(unittest.TestCase): + pass + + +add_function_test( + TestImplicitInitIsMempoolAccessSupported, + "test_is_mempool_access_supported", + test_is_mempool_access_supported, + check_output=False, +) + + +# Peer Access +# ------------------------------------------------------------------------------ + + +def test_is_peer_access_enabled(test, device): + wp.is_peer_access_enabled("cpu", "cpu") + + +class TestImplicitInitIsPeerAccessEnabled(unittest.TestCase): + pass + + +add_function_test( + TestImplicitInitIsPeerAccessEnabled, + "test_is_peer_access_enabled", + test_is_peer_access_enabled, + check_output=False, +) + + +def test_is_peer_access_supported(test, device): + wp.is_peer_access_supported("cpu", "cpu") + + +class TestImplicitInitIsPeerAccessSupported(unittest.TestCase): + pass + + +add_function_test( + TestImplicitInitIsPeerAccessSupported, + "test_is_peer_access_supported", + test_is_peer_access_supported, + check_output=False, +) + + +if __name__ == "__main__": + # Do not clear the kernel cache or call anything that would initialize Warp + # since these tests are specifically aiming to catch issues where Warp isn't + # correctly initialized upon calling certain public APIs. + unittest.main(verbosity=2, failfast=True) From 58e2b8c079cb9b6b28ec3e84879a5c41f83b01e5 Mon Sep 17 00:00:00 2001 From: Christopher Crouzet Date: Wed, 12 Jun 2024 15:30:20 +1200 Subject: [PATCH 4/5] Skip initializing the runtime when running tests Some tests have been introduced for the sole purpose of ensuring that calling some API correctly initializes Warp's runtime if it wasn't initialized already. --- warp/thirdparty/unittest_parallel.py | 3 --- 1 file changed, 3 deletions(-) diff --git a/warp/thirdparty/unittest_parallel.py b/warp/thirdparty/unittest_parallel.py index 79ac956bd..acf7eff0b 100644 --- a/warp/thirdparty/unittest_parallel.py +++ b/warp/thirdparty/unittest_parallel.py @@ -554,9 +554,6 @@ def initialize_test_process(lock, shared_index, args, temp_dir): wp.config.kernel_cache_dir = cache_root_dir wp.build.clear_kernel_cache() - else: - # Initialize Warp is if hasn't been initialized already - wp.init() if __name__ == "__main__": # pragma: no cover From d9f4389f0eb69256c8b714faaed7562c9813bac6 Mon Sep 17 00:00:00 2001 From: Christopher Crouzet Date: Wed, 12 Jun 2024 15:31:50 +1200 Subject: [PATCH 5/5] Add a CI job to test the runtime initialization --- .gitlab-ci.yml | 22 ++++++++++++++++++++++ 1 file changed, 22 insertions(+) diff --git a/.gitlab-ci.yml b/.gitlab-ci.yml index 9a02f5990..c1ecf43e3 100644 --- a/.gitlab-ci.yml +++ b/.gitlab-ci.yml @@ -294,6 +294,28 @@ windows-x86_64 test mgpu: - os/windows - gpu/2x-A5000 + +linux-x86_64 test warp-init: + stage: test + needs: [linux-x86_64 build] + extends: + - .omni_nvks_gpu + - .save_test_report_artifact + before_script: + - echo -e "\\e[0Ksection_start:`date +%s`:install_dependencies[collapsed=true]\\r\\e[0KInstalling dependencies" + - df -h + # Move compiled binaries out of platform-specific directory + - mv warp/bin/linux-x86_64/warp.so warp/bin/ + - mv warp/bin/linux-x86_64/warp-clang.so warp/bin/ + - tools/packman/packman install -l _build/target-deps/python python ${DEFAULT_PYTHON}-linux-x86_64 + - $PYTHON -m venv _venv + - source _venv/bin/activate + - python -m pip install --upgrade pip + - python -m pip install -e . + - echo -e "\\e[0Ksection_end:`date +%s`:install_dependencies\\r\\e[0K" + script: + - python -m warp.tests --junit-report-xml rspec.xml -s autodetect --disable-process-pooling --disable-concurrent-futures --level test -p 'test_implicit_init.py' + # The only purpose of this job is to make sure documentation can be built on Windows. # The output does not get published anywhere, but the website can be viewed in the # artifacts.