Add support for structured types in Device Arrays #7226

testhound · 2021-07-21T16:02:16Z

This pull request addresses lack of support for structured data types with cuda device arrays. The problem is outlined in this issue:

Fixes #6114.

where attempting to set a structured type in a device array fails because there is no check for a DeviceRecord.

gmarkall

I started leaving some comments on the diff, but then I realised that I don't think this works in general. The following:

from numba import cuda
import numpy as np

recwithmat = np.dtype([('i', np.int32),
                       ('j', np.float32, (3, 3))])

recwithrecwithmat = np.dtype([('x', np.int32), ('y', recwithmat)])

arr = np.zeros(1, dtype=recwithrecwithmat)
d_arr = cuda.to_device(arr)

d_arr[0]['y']['i'] = 1
# Prints 0, expected 1
print(d_arr[0]['y']['i'])

d_arr[0]['y']['j'][0, 0] = 2
# Prints 0.0, expected 2.0
print(d_arr[0]['y']['j'][0, 0])


p_arr = cuda.pinned_array(1, dtype=recwithrecwithmat)

p_arr[0]['y']['i'] = 1
# Prints 1 as expected
print(p_arr[0]['y']['i'])

p_arr[0]['y']['j'][0, 0] = 2
# Prints 2.0 as expected
print(p_arr[0]['y']['j'][0, 0])

gives:

whereas I'd expect

under correct operation. Are the get and setitem implementations still incorrect, or am I misunderstanding what this PR fixes?

numba/cuda/api_util.py

numba/cuda/cudadrv/devicearray.py

numba/cuda/tests/cudapy/test_record_dtype.py

gmarkall · 2021-07-27T05:43:02Z

@testhound Is this ready for another review?

testhound · 2021-07-27T06:14:06Z

@testhound Is this ready for another review?

@gmarkall yes it is.

gmarkall

When I run the tests with:

python -m numba.runtests numba.cuda.tests -v -m

I get:

======================================================================
ERROR: test_host_alloc_mapped (numba.cuda.tests.cudadrv.test_host_alloc.TestHostAlloc)
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/gmarkall/numbadev/numba/numba/cuda/tests/cudadrv/test_host_alloc.py", line 43, in test_host_alloc_mapped
    self.assertTrue(all(ary == 123))
  File "/home/gmarkall/numbadev/numba/numba/cuda/cudadrv/devices.py", line 224, in _require_cuda_context
    return fn(*args, **kws)
  File "/home/gmarkall/numbadev/numba/numba/cuda/cudadrv/devicearray.py", line 117, in __getitem__
    return self._do_getitem(item)
AttributeError: 'MappedNDArray' object has no attribute '_do_getitem'

======================================================================
ERROR: test_host_operators (numba.cuda.tests.cudadrv.test_host_alloc.TestHostAlloc)
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/gmarkall/numbadev/numba/numba/cuda/tests/cudadrv/test_host_alloc.py", line 51, in test_host_operators
    ary[:] = range(10)
  File "/home/gmarkall/numbadev/numba/numba/cuda/cudadrv/devices.py", line 224, in _require_cuda_context
    return fn(*args, **kws)
  File "/home/gmarkall/numbadev/numba/numba/cuda/cudadrv/devicearray.py", line 126, in __setitem__
    return self._do_setitem(key, value)
AttributeError: 'MappedNDArray' object has no attribute '_do_setitem'

======================================================================
ERROR: test_managed_array_attach_global (numba.cuda.tests.cudadrv.test_managed_alloc.TestManagedAlloc)
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/gmarkall/numbadev/numba/numba/cuda/tests/cudadrv/test_managed_alloc.py", line 112, in test_managed_array_attach_global
    self._test_managed_array()
  File "/home/gmarkall/numbadev/numba/numba/cuda/tests/cudadrv/test_managed_alloc.py", line 98, in _test_managed_array
    self.assertTrue(all(ary == 123.456))
  File "/home/gmarkall/numbadev/numba/numba/cuda/cudadrv/devices.py", line 224, in _require_cuda_context
    return fn(*args, **kws)
  File "/home/gmarkall/numbadev/numba/numba/cuda/cudadrv/devicearray.py", line 117, in __getitem__
    return self._do_getitem(item)
AttributeError: 'ManagedNDArray' object has no attribute '_do_getitem'

======================================================================
ERROR: test_managed_array_attach_host (numba.cuda.tests.cudadrv.test_managed_alloc.TestManagedAlloc)
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/gmarkall/numbadev/numba/numba/cuda/tests/cudadrv/test_managed_alloc.py", line 115, in test_managed_array_attach_host
    self._test_managed_array()
  File "/home/gmarkall/numbadev/numba/numba/cuda/tests/cudadrv/test_managed_alloc.py", line 98, in _test_managed_array
    self.assertTrue(all(ary == 123.456))
  File "/home/gmarkall/numbadev/numba/numba/cuda/cudadrv/devices.py", line 224, in _require_cuda_context
    return fn(*args, **kws)
  File "/home/gmarkall/numbadev/numba/numba/cuda/cudadrv/devicearray.py", line 117, in __getitem__
    return self._do_getitem(item)
AttributeError: 'ManagedNDArray' object has no attribute '_do_getitem'

======================================================================
FAIL: test_structured_array1 (numba.cuda.tests.cudadrv.test_cuda_devicerecord.TestRecordDtypeWithStructArrays)
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/gmarkall/numbadev/numba/numba/cuda/tests/cudadrv/test_cuda_devicerecord.py", line 147, in test_structured_array1
    self.assertEqual(ary[i]['d'], "%d" % x)
AssertionError: '1\x03\x00\x04' != '1'

----------------------------------------------------------------------

Are these passing on your system?

numba/cuda/cudadrv/devicearray.py

…to their original classes

gmarkall

Many thanks for the changes! This is now looking much closer. I've tested with the simulator and with the hardware target, and all tests pass. I also tested with compute-sanitizer, and this doesn't introduce any new issues. There are some changes I'd suggest:

I think it should be possible to remove the changes to _numba_type_, as suggested on the diff.
The "subarray case" branch in _do_setitem appears to be unused, or untested - does this need removing, or additional tests adding?
In an earlier iteration of the PR, you had a test with a string dtype in the structured type, but it was later removed, I guess because it was failing. I think it was failing because the string was only partially set in the test case. Perhaps it could be reinstated with:

diff --git a/numba/cuda/tests/cudadrv/test_cuda_devicerecord.py b/numba/cuda/tests/cudadrv/test_cuda_devicerecord.py
index d10910446..e2acd34d7 100644
--- a/numba/cuda/tests/cudadrv/test_cuda_devicerecord.py
+++ b/numba/cuda/tests/cudadrv/test_cuda_devicerecord.py
@@ -7,12 +7,14 @@ from numba.cuda.testing import skip_on_cudasim
 from numba.np import numpy_support
 from numba import cuda
 
+N_CHARS = 5
+
 recordtype = np.dtype(
     [
         ('a', np.float64),
         ('b', np.int32),
         ('c', np.complex64),
-        ('d', np.int64)
+        ('d', (np.str_, N_CHARS))
     ],
     align=True
 )
@@ -135,7 +137,7 @@ class TestRecordDtypeWithStructArrays(CUDATestCase):
             ary[i]['a'] = x / 2
             ary[i]['b'] = x
             ary[i]['c'] = x * 1j
-            ary[i]['d'] = x * x
+            ary[i]['d'] = str(x) * N_CHARS
 
     def test_structured_array1(self):
         ary = self.sample1d
@@ -144,7 +146,7 @@ class TestRecordDtypeWithStructArrays(CUDATestCase):
             self.assertEqual(ary[i]['a'], x / 2)
             self.assertEqual(ary[i]['b'], x)
             self.assertEqual(ary[i]['c'], x * 1j)
-            self.assertEqual(ary[i]['d'], x * x)
+            self.assertEqual(ary[i]['d'], str(x) * N_CHARS)
 
     def test_structured_array2(self):
         ary = self.samplerec1darr

numba/cuda/cudadrv/devicearray.py

gmarkall

Many thanks for the update! This tests OK for me locally with the hardware target, compute-sanitizer, and the simulator.

The previous suggestions are addressed, but there are a couple of new suggestions on the new changes - after resolving these I think this will be good from my perspective.

numba/cuda/cudadrv/devicearray.py

gmarkall

Many thanks for the update - this looks good to me!

@stuartarchibald Could this have a buildfarm run please?

testhound · 2021-09-08T03:52:20Z

@gmarkall and @stuartarchibald checking on the status of this PR from three weeks ago.

numba/cuda/cudadrv/devicearray.py

stuartarchibald · 2021-09-08T11:18:56Z

Buildfarm ID: numba_smoketest_cuda_yaml_93.

stuartarchibald · 2021-09-08T11:44:04Z

Buildfarm ID: numba_smoketest_cuda_yaml_93.

This passed on 2dbd81c.

stuartarchibald · 2021-09-15T11:34:45Z

Buildfarm ID: numba_smoketest_cuda_yaml_96.

stuartarchibald · 2021-09-15T14:18:19Z

Buildfarm ID: numba_smoketest_cuda_yaml_96.

Passed.

gmarkall

LGTM with the removed check.

stuartarchibald · 2021-09-15T15:00:13Z

Thanks for the patch @testhound, thanks for reviewing @gmarkall

stuartarchibald · 2021-09-15T15:03:10Z

Presume this closes #6114 or is there more to do with respect to this?

gmarkall · 2021-09-15T15:41:47Z

Presume this closes #6114 or is there more to do with respect to this?

I've updated the description so that merging this PR will close that issue.

stuartarchibald · 2021-09-15T16:09:07Z

Presume this closes #6114 or is there more to do with respect to this?

I've updated the description so that merging this PR will close that issue.

Thanks for confirming and doing the update.

testhound added 4 commits June 22, 2021 14:57

Initial commit for structured device array support with device arrays

2880698

Merge branch 'master' into testhound/set_structured_type_dev_array_6114

c619fc6

Add test cases and fix subarray issues

6b87e61

Merge with upstream master and fix conflicts

237f34e

testhound requested a review from gmarkall as a code owner July 21, 2021 16:02

esc added 3 - Ready for Review CUDA CUDA related issue/PR labels Jul 21, 2021

gmarkall requested changes Jul 22, 2021

View reviewed changes

testhound added 2 commits July 26, 2021 22:28

Fix nested structured arrays and move unit tests

4427f47

Merge branch 'master' into testhound/set_structured_type_dev_array_6114

4e024c4

gmarkall added 4 - Waiting on reviewer Waiting for reviewer to respond to author and removed 3 - Ready for Review labels Jul 27, 2021

gmarkall reviewed Jul 28, 2021

View reviewed changes

numba/cuda/cudadrv/devicearray.py Show resolved Hide resolved

numba/cuda/cudadrv/devicearray.py Show resolved Hide resolved

numba/cuda/cudadrv/devicearray.py Outdated Show resolved Hide resolved

numba/cuda/cudadrv/devicearray.py Outdated Show resolved Hide resolved

testhound added 2 commits July 29, 2021 00:57

Move all variation of getitm and setup out of DeviceNDArrayBase back …

bc6559c

…to their original classes

Merge branch 'master' into testhound/set_structured_type_dev_array_6114

2553c68

gmarkall requested changes Aug 6, 2021

View reviewed changes

numba/cuda/cudadrv/devicearray.py Outdated Show resolved Hide resolved

numba/cuda/cudadrv/devicearray.py Outdated Show resolved Hide resolved

numba/cuda/cudadrv/devicearray.py Outdated Show resolved Hide resolved

gmarkall added 4 - Waiting on author Waiting for author to respond to review and removed 4 - Waiting on reviewer Waiting for reviewer to respond to author labels Aug 11, 2021

testhound added 3 commits August 11, 2021 22:15

Remove unnecessary code and fix test case

661ee7a

Merge branch 'master' into testhound/set_structured_type_dev_array_6114

9377f16

Fix flake8 formatting issue

c83d12d

gmarkall added 4 - Waiting on reviewer Waiting for reviewer to respond to author and removed 4 - Waiting on author Waiting for author to respond to review labels Aug 12, 2021

gmarkall requested changes Aug 12, 2021

View reviewed changes

numba/cuda/cudadrv/devicearray.py Outdated Show resolved Hide resolved

numba/cuda/cudadrv/devicearray.py Outdated Show resolved Hide resolved

gmarkall added 4 - Waiting on author Waiting for author to respond to review and removed 4 - Waiting on reviewer Waiting for reviewer to respond to author labels Aug 12, 2021

Fix comment issues

2dbd81c

gmarkall previously approved these changes Aug 17, 2021

View reviewed changes

gmarkall added 4 - Waiting on CI Review etc done, waiting for CI to finish Pending BuildFarm For PRs that have been reviewed but pending a push through our buildfarm and removed 4 - Waiting on author Waiting for author to respond to review labels Aug 17, 2021

stuartarchibald reviewed Sep 8, 2021

View reviewed changes

numba/cuda/cudadrv/devicearray.py Outdated Show resolved Hide resolved

gmarkall added 4 - Waiting on author Waiting for author to respond to review and removed 4 - Waiting on CI Review etc done, waiting for CI to finish labels Sep 13, 2021

testhound added 2 commits September 13, 2021 12:57

Remove unnecessary error check

2b65b91

Merge branch 'master' into testhound/set_structured_type_dev_array_6114

40a73e8

testhound dismissed gmarkall’s stale review via 40a73e8 September 13, 2021 18:31

stuartarchibald added 4 - Waiting on reviewer Waiting for reviewer to respond to author and removed 4 - Waiting on author Waiting for author to respond to review labels Sep 15, 2021

stuartarchibald added this to the Numba 0.55 RC milestone Sep 15, 2021

stuartarchibald added BuildFarm Passed For PRs that have been through the buildfarm and passed and removed Pending BuildFarm For PRs that have been reviewed but pending a push through our buildfarm labels Sep 15, 2021

gmarkall approved these changes Sep 15, 2021

View reviewed changes

stuartarchibald added 5 - Ready to merge Review and testing done, is ready to merge and removed 4 - Waiting on reviewer Waiting for reviewer to respond to author labels Sep 15, 2021

stuartarchibald approved these changes Sep 15, 2021

View reviewed changes

sklam merged commit 8f3d6e5 into numba:master Sep 16, 2021

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Add support for structured types in Device Arrays #7226

Add support for structured types in Device Arrays #7226

testhound commented Jul 21, 2021 •

edited by gmarkall

gmarkall left a comment

gmarkall commented Jul 27, 2021

testhound commented Jul 27, 2021

gmarkall left a comment

gmarkall left a comment

gmarkall left a comment

gmarkall left a comment

testhound commented Sep 8, 2021

stuartarchibald commented Sep 8, 2021

stuartarchibald commented Sep 8, 2021

stuartarchibald commented Sep 15, 2021

stuartarchibald commented Sep 15, 2021

gmarkall left a comment

stuartarchibald commented Sep 15, 2021

stuartarchibald commented Sep 15, 2021

gmarkall commented Sep 15, 2021

stuartarchibald commented Sep 15, 2021

Add support for structured types in Device Arrays #7226

Add support for structured types in Device Arrays #7226

Conversation

testhound commented Jul 21, 2021 • edited by gmarkall

gmarkall left a comment

Choose a reason for hiding this comment

gmarkall commented Jul 27, 2021

testhound commented Jul 27, 2021

gmarkall left a comment

Choose a reason for hiding this comment

gmarkall left a comment

Choose a reason for hiding this comment

gmarkall left a comment

Choose a reason for hiding this comment

gmarkall left a comment

Choose a reason for hiding this comment

testhound commented Sep 8, 2021

stuartarchibald commented Sep 8, 2021

stuartarchibald commented Sep 8, 2021

stuartarchibald commented Sep 15, 2021

stuartarchibald commented Sep 15, 2021

gmarkall left a comment

Choose a reason for hiding this comment

stuartarchibald commented Sep 15, 2021

stuartarchibald commented Sep 15, 2021

gmarkall commented Sep 15, 2021

stuartarchibald commented Sep 15, 2021

testhound commented Jul 21, 2021 •

edited by gmarkall