diff --git a/cuda_core/cuda/core/experimental/_program.py b/cuda_core/cuda/core/experimental/_program.py index ae5928ee..ec0778a3 100644 --- a/cuda_core/cuda/core/experimental/_program.py +++ b/cuda_core/cuda/core/experimental/_program.py @@ -14,9 +14,9 @@ class Program: _supported_target_type = ("ptx", "cubin", "ltoir", ) def __init__(self, code, code_type): + self._handle = None if code_type not in self._supported_code_type: raise NotImplementedError - self._handle = None if code_type.lower() == "c++": if not isinstance(code, str): diff --git a/cuda_core/tests/conftest.py b/cuda_core/tests/conftest.py new file mode 100644 index 00000000..3ff6ce08 --- /dev/null +++ b/cuda_core/tests/conftest.py @@ -0,0 +1,16 @@ +# Copyright 2024 NVIDIA Corporation. All rights reserved. +# +# Please refer to the NVIDIA end user license agreement (EULA) associated +# with this source code for terms and conditions that govern your use of +# this software. Any use, reproduction, disclosure, or distribution of +# this software and related documentation outside the terms of the EULA +# is strictly prohibited. + +from cuda.core.experimental._device import Device +import pytest + +@pytest.fixture(scope="module") +def init_cuda(): + device = Device() + device.set_current() + \ No newline at end of file diff --git a/cuda_core/tests/example_tests/__init__.py b/cuda_core/tests/example_tests/__init__.py new file mode 100644 index 00000000..e69de29b diff --git a/cuda_core/tests/example_tests/test_basic_examples.py b/cuda_core/tests/example_tests/test_basic_examples.py new file mode 100644 index 00000000..e490892d --- /dev/null +++ b/cuda_core/tests/example_tests/test_basic_examples.py @@ -0,0 +1,25 @@ +# Copyright 2024 NVIDIA Corporation. All rights reserved. +# +# Please refer to the NVIDIA end user license agreement (EULA) associated +# with this source code for terms and conditions that govern your use of +# this software. Any use, reproduction, disclosure, or distribution of +# this software and related documentation outside the terms of the EULA +# is strictly prohibited. + +# If we have subcategories of examples in the future, this file can be split along those lines + +from .utils import run_example +import os +import glob +import pytest + +samples_path = os.path.join( + os.path.dirname(__file__), '..', '..', 'examples') +sample_files = glob.glob(samples_path+'**/*.py', recursive=True) +@pytest.mark.parametrize( + 'example', sample_files +) +class TestExamples: + def test_example(self, example): + filename = os.path.basename(example) + run_example(samples_path, example) diff --git a/cuda_core/tests/example_tests/utils.py b/cuda_core/tests/example_tests/utils.py new file mode 100644 index 00000000..5f4e14b0 --- /dev/null +++ b/cuda_core/tests/example_tests/utils.py @@ -0,0 +1,54 @@ +# Copyright 2024 NVIDIA Corporation. All rights reserved. +# +# Please refer to the NVIDIA end user license agreement (EULA) associated +# with this source code for terms and conditions that govern your use of +# this software. Any use, reproduction, disclosure, or distribution of +# this software and related documentation outside the terms of the EULA +# is strictly prohibited. + +from cuda import cuda +import gc +import os +import sys +import pytest +import cupy as cp + +class SampleTestError(Exception): + pass + +def parse_python_script(filepath): + if not filepath.endswith('.py'): + raise ValueError(f"{filepath} not supported") + with open(filepath, "r", encoding='utf-8') as f: + script = f.read() + return script + + +def run_example(samples_path, filename, env=None): + fullpath = os.path.join(samples_path, filename) + script = parse_python_script(fullpath) + try: + old_argv = sys.argv + sys.argv = [fullpath] + old_sys_path = sys.path.copy() + sys.path.append(samples_path) + exec(script, env if env else {}) + except ImportError as e: + # for samples requiring any of optional dependencies + for m in ('cupy',): + if f"No module named '{m}'" in str(e): + pytest.skip(f'{m} not installed, skipping related tests') + break + else: + raise + except Exception as e: + msg = "\n" + msg += f'Got error ({filename}):\n' + msg += str(e) + raise SampleTestError(msg) from e + finally: + sys.path = old_sys_path + sys.argv = old_argv + # further reduce the memory watermark + gc.collect() + cp.get_default_memory_pool().free_all_blocks() diff --git a/cuda_core/tests/test_device.py b/cuda_core/tests/test_device.py new file mode 100644 index 00000000..653dac06 --- /dev/null +++ b/cuda_core/tests/test_device.py @@ -0,0 +1,67 @@ +# Copyright 2024 NVIDIA Corporation. All rights reserved. +# +# Please refer to the NVIDIA end user license agreement (EULA) associated +# with this source code for terms and conditions that govern your use of +# this software. Any use, reproduction, disclosure, or distribution of +# this software and related documentation outside the terms of the EULA +# is strictly prohibited. + +from cuda import cuda, cudart +from cuda.core.experimental._device import Device +from cuda.core.experimental._utils import handle_return, ComputeCapability, CUDAError, \ + precondition +import pytest + +def test_device_repr(): + device = Device(0) + assert str(device).startswith('= 11040: + uuid = handle_return(cuda.cuDeviceGetUuid_v2(device.device_id)) + else: + uuid = handle_return(cuda.cuDeviceGetUuid(device.device_id)) + uuid = uuid.bytes.hex() + expected_uuid = f"{uuid[:8]}-{uuid[8:12]}-{uuid[12:16]}-{uuid[16:20]}-{uuid[20:]}" + assert device.uuid == expected_uuid + +def test_name(): + device = Device() + name = handle_return(cuda.cuDeviceGetName(128, device.device_id)) + name = name.split(b'\0')[0] + assert device.name == name.decode() + +def test_compute_capability(): + device = Device() + major = handle_return(cudart.cudaDeviceGetAttribute( + cudart.cudaDeviceAttr.cudaDevAttrComputeCapabilityMajor, device.device_id)) + minor = handle_return(cudart.cudaDeviceGetAttribute( + cudart.cudaDeviceAttr.cudaDevAttrComputeCapabilityMinor, device.device_id)) + expected_cc = ComputeCapability(major, minor) + assert device.compute_capability == expected_cc + \ No newline at end of file diff --git a/cuda_core/tests/test_event.py b/cuda_core/tests/test_event.py new file mode 100644 index 00000000..b6cfe647 --- /dev/null +++ b/cuda_core/tests/test_event.py @@ -0,0 +1,39 @@ +# Copyright 2024 NVIDIA Corporation. All rights reserved. +# +# Please refer to the NVIDIA end user license agreement (EULA) associated +# with this source code for terms and conditions that govern your use of +# this software. Any use, reproduction, disclosure, or distribution of +# this software and related documentation outside the terms of the EULA +# is strictly prohibited. + +from cuda import cuda +from cuda.core.experimental._event import EventOptions, Event +from cuda.core.experimental._utils import handle_return +from cuda.core.experimental._device import Device +import pytest + +def test_is_timing_disabled(): + options = EventOptions(enable_timing=False) + event = Event._init(options) + assert event.is_timing_disabled == True + +def test_is_sync_busy_waited(): + options = EventOptions(busy_waited_sync=True) + event = Event._init(options) + assert event.is_sync_busy_waited == True + +def test_sync(): + options = EventOptions() + event = Event._init(options) + event.sync() + assert event.is_done == True + +def test_is_done(): + options = EventOptions() + event = Event._init(options) + assert event.is_done == True + +def test_handle(): + options = EventOptions() + event = Event._init(options) + assert isinstance(event.handle, int) diff --git a/cuda_core/tests/test_launcher.py b/cuda_core/tests/test_launcher.py new file mode 100644 index 00000000..92dfc726 --- /dev/null +++ b/cuda_core/tests/test_launcher.py @@ -0,0 +1,66 @@ +# Copyright 2024 NVIDIA Corporation. All rights reserved. +# +# Please refer to the NVIDIA end user license agreement (EULA) associated +# with this source code for terms and conditions that govern your use of +# this software. Any use, reproduction, disclosure, or distribution of +# this software and related documentation outside the terms of the EULA +# is strictly prohibited. + +from cuda import cuda +from cuda.core.experimental._launcher import LaunchConfig +from cuda.core.experimental._stream import Stream +from cuda.core.experimental._device import Device +from cuda.core.experimental._utils import handle_return +import pytest + +def test_launch_config_init(): + config = LaunchConfig(grid=(1, 1, 1), block=(1, 1, 1), stream=None, shmem_size=0) + assert config.grid == (1, 1, 1) + assert config.block == (1, 1, 1) + assert config.stream is None + assert config.shmem_size == 0 + + config = LaunchConfig(grid=(2, 2, 2), block=(2, 2, 2), stream=Device().create_stream(), shmem_size=1024) + assert config.grid == (2, 2, 2) + assert config.block == (2, 2, 2) + assert isinstance(config.stream, Stream) + assert config.shmem_size == 1024 + +def test_launch_config_cast_to_3_tuple(): + config = LaunchConfig(grid=1, block=1) + assert config._cast_to_3_tuple(1) == (1, 1, 1) + assert config._cast_to_3_tuple((1, 2)) == (1, 2, 1) + assert config._cast_to_3_tuple((1, 2, 3)) == (1, 2, 3) + + # Edge cases + assert config._cast_to_3_tuple(999) == (999, 1, 1) + assert config._cast_to_3_tuple((999, 888)) == (999, 888, 1) + assert config._cast_to_3_tuple((999, 888, 777)) == (999, 888, 777) + +def test_launch_config_invalid_values(): + with pytest.raises(ValueError): + LaunchConfig(grid=0, block=1) + + with pytest.raises(ValueError): + LaunchConfig(grid=(0, 1), block=1) + + with pytest.raises(ValueError): + LaunchConfig(grid=(1, 1, 1), block=0) + + with pytest.raises(ValueError): + LaunchConfig(grid=(1, 1, 1), block=(0, 1)) + +def test_launch_config_stream(): + stream = Device().create_stream() + config = LaunchConfig(grid=(1, 1, 1), block=(1, 1, 1), stream=stream, shmem_size=0) + assert config.stream == stream + + with pytest.raises(ValueError): + LaunchConfig(grid=(1, 1, 1), block=(1, 1, 1), stream="invalid_stream", shmem_size=0) + +def test_launch_config_shmem_size(): + config = LaunchConfig(grid=(1, 1, 1), block=(1, 1, 1), stream=None, shmem_size=2048) + assert config.shmem_size == 2048 + + config = LaunchConfig(grid=(1, 1, 1), block=(1, 1, 1), stream=None) + assert config.shmem_size == 0 diff --git a/cuda_core/tests/test_memory.py b/cuda_core/tests/test_memory.py new file mode 100644 index 00000000..40855268 --- /dev/null +++ b/cuda_core/tests/test_memory.py @@ -0,0 +1,197 @@ +# Copyright 2024 NVIDIA Corporation. All rights reserved. +# +# Please refer to the NVIDIA end user license agreement (EULA) associated +# with this source code for terms and conditions that govern your use of +# this software. Any use, reproduction, disclosure, or distribution of +# this software and related documentation outside the terms of the EULA +# is strictly prohibited. + +from cuda import cuda +from cuda.core.experimental._memory import Buffer, MemoryResource +from cuda.core.experimental._device import Device +from cuda.core.experimental._utils import handle_return +import ctypes +import pytest + +class DummyDeviceMemoryResource(MemoryResource): + def __init__(self, device): + self.device = device + + def allocate(self, size, stream=None) -> Buffer: + ptr = handle_return(cuda.cuMemAlloc(size)) + return Buffer(ptr=ptr, size=size, mr=self) + + def deallocate(self, ptr, size, stream=None): + handle_return(cuda.cuMemFree(ptr)) + + @property + def is_device_accessible(self) -> bool: + return True + + @property + def is_host_accessible(self) -> bool: + return False + + @property + def device_id(self) -> int: + return 0 + +class DummyHostMemoryResource(MemoryResource): + def __init__(self): + pass + + def allocate(self, size, stream=None) -> Buffer: + # Allocate a ctypes buffer of size `size` + ptr = (ctypes.c_byte * size)() + return Buffer(ptr=ptr, size=size, mr=self) + + def deallocate(self, ptr, size, stream=None): + #the memory is deallocated per the ctypes deallocation at garbage collection time + pass + + @property + def is_device_accessible(self) -> bool: + return False + + @property + def is_host_accessible(self) -> bool: + return True + + @property + def device_id(self) -> int: + raise RuntimeError("the pinned memory resource is not bound to any GPU") + +class DummyUnifiedMemoryResource(MemoryResource): + def __init__(self, device): + self.device = device + + def allocate(self, size, stream=None) -> Buffer: + ptr = handle_return(cuda.cuMemAllocManaged(size, cuda.CUmemAttach_flags.CU_MEM_ATTACH_GLOBAL.value)) + return Buffer(ptr=ptr, size=size, mr=self) + + def deallocate(self, ptr, size, stream=None): + handle_return(cuda.cuMemFree(ptr)) + + @property + def is_device_accessible(self) -> bool: + return True + + @property + def is_host_accessible(self) -> bool: + return True + + @property + def device_id(self) -> int: + return 0 + +class DummyPinnedMemoryResource(MemoryResource): + def __init__(self, device): + self.device = device + + def allocate(self, size, stream=None) -> Buffer: + ptr = handle_return(cuda.cuMemAllocHost(size)) + return Buffer(ptr=ptr, size=size, mr=self) + + def deallocate(self, ptr, size, stream=None): + handle_return(cuda.cuMemFreeHost(ptr)) + + @property + def is_device_accessible(self) -> bool: + return True + + @property + def is_host_accessible(self) -> bool: + return True + + @property + def device_id(self) -> int: + raise RuntimeError("the pinned memory resource is not bound to any GPU") + +def buffer_initialization(dummy_mr : MemoryResource): + buffer = dummy_mr.allocate(size=1024) + assert buffer.handle != 0 + assert buffer.size == 1024 + assert buffer.memory_resource == dummy_mr + assert buffer.is_device_accessible == dummy_mr.is_device_accessible + assert buffer.is_host_accessible == dummy_mr.is_host_accessible + buffer.close() + +def test_buffer_initialization(): + device = Device() + device.set_current() + buffer_initialization(DummyDeviceMemoryResource(device)) + buffer_initialization(DummyHostMemoryResource()) + buffer_initialization(DummyUnifiedMemoryResource(device)) + buffer_initialization(DummyPinnedMemoryResource(device)) + +def buffer_copy_to(dummy_mr : MemoryResource, device : Device, check = False): + src_buffer = dummy_mr.allocate(size=1024) + dst_buffer = dummy_mr.allocate(size=1024) + stream = device.create_stream() + + if check: + src_ptr = ctypes.cast(src_buffer.handle, ctypes.POINTER(ctypes.c_byte)) + for i in range(1024): + src_ptr[i] = ctypes.c_byte(i) + + src_buffer.copy_to(dst_buffer, stream=stream) + device.sync() + + if check: + dst_ptr = ctypes.cast(dst_buffer.handle, ctypes.POINTER(ctypes.c_byte)) + + for i in range(10): + assert dst_ptr[i] == src_ptr[i] + + dst_buffer.close() + src_buffer.close() + +def test_buffer_copy_to(): + device = Device() + device.set_current() + buffer_copy_to(DummyDeviceMemoryResource(device), device) + buffer_copy_to(DummyUnifiedMemoryResource(device), device) + buffer_copy_to(DummyPinnedMemoryResource(device), device, check = True) + +def buffer_copy_from(dummy_mr : MemoryResource, device, check = False): + src_buffer = dummy_mr.allocate(size=1024) + dst_buffer = dummy_mr.allocate(size=1024) + stream = device.create_stream() + + if check: + src_ptr = ctypes.cast(src_buffer.handle, ctypes.POINTER(ctypes.c_byte)) + for i in range(1024): + src_ptr[i] = ctypes.c_byte(i) + + dst_buffer.copy_from(src_buffer, stream=stream) + device.sync() + + if check: + dst_ptr = ctypes.cast(dst_buffer.handle, ctypes.POINTER(ctypes.c_byte)) + + for i in range(10): + assert dst_ptr[i] == src_ptr[i] + + dst_buffer.close() + src_buffer.close() + +def test_buffer_copy_from(): + device = Device() + device.set_current() + buffer_copy_from(DummyDeviceMemoryResource(device), device) + buffer_copy_from(DummyUnifiedMemoryResource(device), device) + buffer_copy_from(DummyPinnedMemoryResource(device), device, check = True) + +def buffer_close(dummy_mr : MemoryResource): + buffer = dummy_mr.allocate(size=1024) + buffer.close() + assert buffer.handle == 0 + assert buffer.memory_resource == None + +def test_buffer_close(): + device = Device() + device.set_current() + buffer_close(DummyDeviceMemoryResource(device)) + buffer_close(DummyHostMemoryResource()) + buffer_close(DummyUnifiedMemoryResource(device)) + buffer_close(DummyPinnedMemoryResource(device)) diff --git a/cuda_core/tests/test_module.py b/cuda_core/tests/test_module.py new file mode 100644 index 00000000..cc5cf57c --- /dev/null +++ b/cuda_core/tests/test_module.py @@ -0,0 +1,39 @@ +# Copyright 2024 NVIDIA Corporation. All rights reserved. +# +# Please refer to the NVIDIA end user license agreement (EULA) associated +# with this source code for terms and conditions that govern your use of +# this software. Any use, reproduction, disclosure, or distribution of +# this software and related documentation outside the terms of the EULA +# is strictly prohibited. + +from cuda import cuda +from cuda.core.experimental._device import Device +from cuda.core.experimental._module import Kernel, ObjectCode +from cuda.core.experimental._utils import handle_return +import pytest + +def test_object_code_initialization(): + # Test with supported code types + for code_type in ["cubin", "ptx", "fatbin"]: + module_data = b"dummy_data" + obj_code = ObjectCode(module_data, code_type) + assert obj_code._code_type == code_type + assert obj_code._module == module_data + assert obj_code._handle is not None + + # Test with unsupported code type + with pytest.raises(ValueError): + ObjectCode(b"dummy_data", "unsupported_code_type") + +#TODO add ObjectCode tests which provide the appropriate data for cuLibraryLoadFromFile +def test_object_code_initialization_with_str(): + assert True + +def test_object_code_initialization_with_jit_options(): + assert True + +def test_object_code_get_kernel(): + assert True + +def test_kernel_from_obj(): + assert True diff --git a/cuda_core/tests/test_program.py b/cuda_core/tests/test_program.py new file mode 100644 index 00000000..39ce4dc6 --- /dev/null +++ b/cuda_core/tests/test_program.py @@ -0,0 +1,58 @@ +# Copyright 2024 NVIDIA Corporation. All rights reserved. +# +# Please refer to the NVIDIA end user license agreement (EULA) associated +# with this source code for terms and conditions that govern your use of +# this software. Any use, reproduction, disclosure, or distribution of +# this software and related documentation outside the terms of the EULA +# is strictly prohibited. + +from cuda.core.experimental._program import Program +from cuda.core.experimental._module import ObjectCode, Kernel +from cuda.core.experimental._device import Device +import pytest + +def test_program_init_valid_code_type(): + code = "extern \"C\" __global__ void my_kernel() {}" + program = Program(code, "c++") + assert program.backend == "nvrtc" + assert program.handle is not None + +def test_program_init_invalid_code_type(): + code = "extern \"C\" __global__ void my_kernel() {}" + with pytest.raises(NotImplementedError): + Program(code, "python") + +def test_program_init_invalid_code_format(): + code = 12345 + with pytest.raises(TypeError): + Program(code, "c++") + +def test_program_compile_valid_target_type(): + code = "extern \"C\" __global__ void my_kernel() {}" + program = Program(code, "c++") + object_code = program.compile("ptx") + kernel = object_code.get_kernel("my_kernel") + assert isinstance(object_code, ObjectCode) + assert isinstance(kernel, Kernel) + +def test_program_compile_invalid_target_type(): + code = "extern \"C\" __global__ void my_kernel() {}" + program = Program(code, "c++") + with pytest.raises(NotImplementedError): + program.compile("invalid_target") + +def test_program_backend_property(): + code = "extern \"C\" __global__ void my_kernel() {}" + program = Program(code, "c++") + assert program.backend == "nvrtc" + +def test_program_handle_property(): + code = "extern \"C\" __global__ void my_kernel() {}" + program = Program(code, "c++") + assert program.handle is not None + +def test_program_close(): + code = "extern \"C\" __global__ void my_kernel() {}" + program = Program(code, "c++") + program.close() + assert program.handle is None diff --git a/cuda_core/tests/test_stream.py b/cuda_core/tests/test_stream.py new file mode 100644 index 00000000..e0a98c18 --- /dev/null +++ b/cuda_core/tests/test_stream.py @@ -0,0 +1,86 @@ +# Copyright 2024 NVIDIA Corporation. All rights reserved. +# +# Please refer to the NVIDIA end user license agreement (EULA) associated +# with this source code for terms and conditions that govern your use of +# this software. Any use, reproduction, disclosure, or distribution of +# this software and related documentation outside the terms of the EULA +# is strictly prohibited. + +from cuda.core.experimental._stream import Stream, StreamOptions, LEGACY_DEFAULT_STREAM, PER_THREAD_DEFAULT_STREAM, default_stream +from cuda.core.experimental._event import Event, EventOptions +from cuda.core.experimental._device import Device +import pytest + +def test_stream_init(): + with pytest.raises(NotImplementedError): + Stream() + +def test_stream_init_with_options(): + stream = Stream._init(options=StreamOptions(nonblocking=True, priority=0)) + assert stream.is_nonblocking is True + assert stream.priority == 0 + +def test_stream_handle(): + stream = Stream._init(options=StreamOptions()) + assert isinstance(stream.handle, int) + +def test_stream_is_nonblocking(): + stream = Stream._init(options=StreamOptions(nonblocking=True)) + assert stream.is_nonblocking is True + +def test_stream_priority(): + stream = Stream._init(options=StreamOptions(priority=0)) + assert stream.priority == 0 + stream = Stream._init(options=StreamOptions(priority=-1)) + assert stream.priority == -1 + with pytest.raises(ValueError): + stream = Stream._init(options=StreamOptions(priority=1)) + +def test_stream_sync(): + stream = Stream._init(options=StreamOptions()) + stream.sync() # Should not raise any exceptions + +def test_stream_record(): + stream = Stream._init(options=StreamOptions()) + event = stream.record() + assert isinstance(event, Event) + +def test_stream_record_invalid_event(): + stream = Stream._init(options=StreamOptions()) + with pytest.raises(TypeError): + stream.record(event="invalid_event") + +def test_stream_wait_event(): + stream = Stream._init(options=StreamOptions()) + event = Event._init() + stream.record(event) + stream.wait(event) # Should not raise any exceptions + +def test_stream_wait_invalid_event(): + stream = Stream._init(options=StreamOptions()) + with pytest.raises(ValueError): + stream.wait(event_or_stream="invalid_event") + +def test_stream_device(): + stream = Stream._init(options=StreamOptions()) + device = stream.device + assert isinstance(device, Device) + +def test_stream_context(): + stream = Stream._init(options=StreamOptions()) + context = stream.context + assert context is not None + +def test_stream_from_handle(): + stream = Stream.from_handle(0) + assert isinstance(stream, Stream) + +def test_legacy_default_stream(): + assert isinstance(LEGACY_DEFAULT_STREAM, Stream) + +def test_per_thread_default_stream(): + assert isinstance(PER_THREAD_DEFAULT_STREAM, Stream) + +def test_default_stream(): + stream = default_stream() + assert isinstance(stream, Stream)