def test_eq_hip(self):
null0 = self.stream
null1 = cuda.Stream(True)
null2 = cuda.Stream(True)
null3 = cuda.Stream()
assert null0 == null1
assert null1 == null2
assert null2 != null3
def test_eq(self):
null0 = cuda.Stream.null
null1 = cuda.Stream(True)
null2 = cuda.Stream(True)
null3 = cuda.Stream()
assert null0 == null1
assert null1 == null2
assert null2 != null3
def test_eq(self):
null0 = cuda.Stream.null
null1 = cuda.Stream(True)
null2 = cuda.Stream(True)
null3 = cuda.Stream()
self.assertEqual(null0, null1)
self.assertEqual(null1, null2)
self.assertNotEqual(null2, null3)
def test_with_statement(self):
stream1 = cuda.Stream()
stream2 = cuda.Stream()
self.assertEqual(cuda.Stream.null, cuda.get_current_stream())
with stream1:
self.assertEqual(stream1, cuda.get_current_stream())
with stream2:
self.assertEqual(stream2, cuda.get_current_stream())
self.assertEqual(stream1, cuda.get_current_stream())
self.assertEqual(cuda.Stream.null, cuda.get_current_stream())
def test_with_statement(self):
stream1 = cuda.Stream()
stream2 = cuda.Stream()
assert cuda.Stream.null == cuda.get_current_stream()
with stream1:
assert stream1 == cuda.get_current_stream()
with stream2:
assert stream2 == cuda.get_current_stream()
assert stream1 == cuda.get_current_stream()
assert cuda.Stream.null == cuda.get_current_stream()
def test_with_statement(self):
stream1 = cuda.Stream()
stream2 = cuda.Stream()
assert self.stream == cuda.get_current_stream()
with stream1:
assert stream1 == cuda.get_current_stream()
with stream2:
assert stream2 == cuda.get_current_stream()
assert stream1 == cuda.get_current_stream()
# self.stream is "forgotten"!
assert cuda.Stream.null == cuda.get_current_stream()
def _get_stream(self, stream_name):
if stream_name == 'null':
return cuda.Stream.null
elif stream_name == 'ptds':
return cuda.Stream.ptds
else:
return cuda.Stream()
def test_mix_use_context(self):
# See cupy/cupy#5143
s1 = cuda.Stream()
s2 = cuda.Stream()
s3 = cuda.Stream()
assert cuda.get_current_stream() == self.stream
with s1:
assert cuda.get_current_stream() == s1
s2.use()
assert cuda.get_current_stream() == s2
with s3:
assert cuda.get_current_stream() == s3
del s2
assert cuda.get_current_stream() == s1
# self.stream is "forgotten"!
assert cuda.get_current_stream() == cuda.Stream.null
def setUp(self):
if self.stream == 'null':
self.stream = cuda.Stream.null
elif self.stream == 'new':
self.stream = cuda.Stream()
self.old_ver = _util.CUDA_ARRAY_INTERFACE_EXPORT_VERSION
_util.CUDA_ARRAY_INTERFACE_EXPORT_VERSION = self.ver
def test_per_device_failure(self):
with cuda.Device(0):
stream0 = cuda.Stream()
with cuda.Device(1):
with pytest.raises(RuntimeError):
with stream0:
pass
with pytest.raises(RuntimeError):
stream0.use()
def test_del(self):
stream = cuda.Stream().use()
stream_ptr = stream.ptr
x = from_data.array([1, 2, 3])
del stream
# Want to test cudaStreamDestory is issued, but
# runtime.streamQuery(stream_ptr) causes SEGV. We cannot test...
del stream_ptr
del x
def test_per_device(self):
with cuda.Device(0):
stream0 = cuda.Stream()
with stream0:
assert stream0 == cuda.get_current_stream()
with cuda.Device(1):
assert stream0 != cuda.get_current_stream()
assert cuda.Stream.null == cuda.get_current_stream()
assert stream0 == cuda.get_current_stream()
def check_del(self, null):
stream = cuda.Stream(null=null).use()
stream_ptr = stream.ptr
x = from_data.array([1, 2, 3])
del stream
assert cuda.Stream.null == cuda.get_current_stream()
# Want to test cudaStreamDestory is issued, but
# runtime.streamQuery(stream_ptr) causes SEGV. We cannot test...
del stream_ptr
del x
def test_del(self):
stream = cuda.Stream().use()
stream_ptr = stream.ptr
x = from_data.array([1, 2, 3])
del stream
self.assertEqual(cuda.Stream.null, cuda.get_current_stream())
# Want to test cudaStreamDestory is issued, but
# runtime.streamQuery(stream_ptr) causes SEGV. We cannot test...
del stream_ptr
del x
def test_copy_multi_device_with_stream(self):
# Kernel that takes long enough then finally writes values.
kern = cupy.RawKernel(
_test_copy_multi_device_with_stream_src, 'wait_and_write')
# Allocates a memory and launches the kernel on a device with its
# stream.
with cuda.Device(0):
with cuda.Stream():
a = cupy.zeros((2,), dtype=numpy.uint64)
kern((1,), (1,), a)
# D2D copy to another device with another stream should get the
# original values of the memory before the kernel on the first device
# finally makes the write.
with cuda.Device(1):
with cuda.Stream():
b = a.copy()
testing.assert_array_equal(
b, numpy.array([0, 0], dtype=numpy.uint64))
def test_copy_multi_device_with_stream(self):
# Kernel that takes long enough then finally writes values.
src = _test_copy_multi_device_with_stream_src
if runtime.is_hip and driver.get_build_version() >= 5_00_00000:
src = '#include <ctime>\n' + src
kern = cupy.RawKernel(src, 'wait_and_write')
# Allocates a memory and launches the kernel on a device with its
# stream.
with cuda.Device(0):
# Keep this stream alive over the D2D copy below for HIP
with cuda.Stream() as s1: # NOQA
a = cupy.zeros((2,), dtype=numpy.uint64)
kern((1,), (1,), a)
# D2D copy to another device with another stream should get the
# original values of the memory before the kernel on the first device
# finally makes the write.
with cuda.Device(1):
with cuda.Stream():
b = a.copy()
testing.assert_array_equal(
b, numpy.array([0, 0], dtype=numpy.uint64))
def f1(barrier, errors):
global s1
tid = barrier.wait()
try:
s1 = cuda.Stream()
barrier.wait() # until t2 starts
s1.use()
barrier.wait() # until t2 uses the stream
s1 = None
gc.collect()
barrier.wait() # until t2 decrefs the stream
assert cuda.get_current_stream() is not None
cupy.arange(10)
errors[tid] = False
except Exception as e:
print(f'error in {tid}: {e}')
def test_eq_cuda(self):
null0 = self.stream
if self.stream == cuda.Stream.null:
null1 = cuda.Stream(True)
null2 = cuda.Stream(True)
null3 = cuda.Stream(ptds=True)
else:
null1 = cuda.Stream(ptds=True)
null2 = cuda.Stream(ptds=True)
null3 = cuda.Stream(True)
null4 = cuda.Stream()
assert null0 == null1
assert null1 == null2
assert null2 != null3
assert null2 != null4
def test_get_and_add_callback(self):
N = 100
cupy_arrays = [testing.shaped_random((2, 3)) for _ in range(N)]
if not cuda.runtime.is_hip:
stream = cuda.Stream.null
else:
# adding callbacks to the null stream in HIP would segfault...
stream = cuda.Stream()
out = []
for i in range(N):
numpy_array = cupy_arrays[i].get(stream=stream)
stream.add_callback(lambda _, __, t: out.append(t[0]),
(i, numpy_array))
stream.synchronize()
assert out == list(range(N))
def test_get_and_add_callback(self):
N = 100
cupy_arrays = [testing.shaped_random((2, 3)) for _ in range(N)]
if not cuda.runtime.is_hip:
stream = self.stream
else:
# adding callbacks to the null stream in HIP would segfault...
stream = cuda.Stream()
out = []
stream_list = []
def _callback(s, _, t):
out.append(t[0])
stream_list.append(s.ptr)
for i in range(N):
numpy_array = cupy_arrays[i].get(stream=stream)
stream.add_callback(_callback, (i, numpy_array))
stream.synchronize()
assert out == list(range(N))
assert all(s == stream.ptr for s in stream_list)
def test_use(self):
stream1 = cuda.Stream().use()
assert stream1 == cuda.get_current_stream()
self.stream.use()
assert self.stream == cuda.get_current_stream()
def test_use(self):
stream1 = cuda.Stream().use()
assert stream1 == cuda.get_current_stream()
cuda.Stream.null.use()
assert cuda.Stream.null == cuda.get_current_stream()
def _init(self):
if not self.initted:
self.start_gpu = cuda.Event()
self.stop_gpu = cuda.Event()
self.stream = cuda.Stream()
self.initted = True
def test_hash(self):
hash(self.stream)
hash(cuda.Stream(True))
hash(cuda.Stream(False))
mapping = {cuda.Stream(): 1, cuda.Stream(): 2} # noqa
def setUp(self):
self.stream = cuda.Stream(null=True)
def test_use(self):
stream1 = cuda.Stream().use()
self.assertEqual(stream1, cuda.get_current_stream())
cuda.Stream.null.use()
self.assertEqual(cuda.Stream.null, cuda.get_current_stream())
def set_stream(self, stream=None):
if stream is None:
stream = cuda.Stream()
curand.setStream(self._generator, stream.ptr)
def setUp(self):
self.stream = cuda.Stream()
