Update CtypesLinker::add_cu error message to include fp16 usage by testhound · Pull Request #8744 · numba/numba

testhound · 2023-02-08T07:10:32Z

This pull request update the error message on the add_cu method in CtypesLinker to include fp16. The reason for the update is CUDA fp16 support requires linking of cu files in it's internal implementation.

testhound · 2023-02-08T16:09:11Z

gpuci run tests

testhound · 2023-02-13T16:51:43Z

gpuci run tests

gmarkall · 2023-02-13T16:56:39Z

gpuci run tests

gmarkall

Thanks for the PR. In addition to the changed test, I think it would be wise to have a test that causes the exception to be raised through the use of an unsupported operation on float16 data.

A couple of other questions:

Is it possible to identify what the user is doing (either using float16 or linking in additional .cu files) and provide them an error message that references specifically what they're doing instead of providing a catch-all generic message?
Is the documentation of this restriction sufficient, or are some additions needed in the docs?

gmarkall · 2023-02-14T12:35:23Z

+        raise NotImplementedError("Support for fp16 or linking CUDA source "
+                                  "files requires the use of the Nvidia CUDA "
+                                  "bindings and enabling the "
+                                  "NUMBA_CUDA_USE_NVIDIA_BINDING environment "
+                                  "variable. See documentation for me details.")


Suggested changes:

fp16 should be float16 -I noticed that the code in Numba is unusual in Python referring to 16-bit floating point values as fp16 rather than float16 which seems to be used everywhere else - for example, the NumPy dtype. This is an inconsistency I should probably have realised much sooner, so I'd like to try not introducing more instances of it.

Nvidia is usually written NVIDIA

Instead of saying the environment variable should be enabled, say that it should be set to 1.

See documentation for me details needs the grammar correcting, and it would be helpful to provide a link to the relevant documentation. Where is the user supposed to look?

testhound · 2023-02-15T02:36:39Z

gpuci run tests

testhound · 2023-02-15T02:57:04Z

@gmarkall error message updated.

gmarkall · 2023-02-15T10:22:25Z

Thanks. What do you think about the other two questions from the review?

testhound · 2023-02-15T15:54:01Z

Thanks. What do you think about the other two questions from the review?

@gmarkall I don't think it's possible to determine if this is a float 16 issue or the user is linking in .cu files; in both cases we end up in the CtypesLinker::add_cu which throws the exception. We would need to catch this higher up in the stack somewhere.

I think the clarified error message along with the link to the documentation is sufficient.

gmarkall · 2023-02-15T17:33:15Z

I don't think it's possible to determine if this is a float 16 issue or the user is linking in .cu files; in both cases we end up in the CtypesLinker::add_cu which throws the exception. We would need to catch this higher up in the stack somewhere.

Is there anything wrong with this approach:

diff --git a/numba/cuda/dispatcher.py b/numba/cuda/dispatcher.py
index a79120d85..bb85093f9 100644
--- a/numba/cuda/dispatcher.py
+++ b/numba/cuda/dispatcher.py
@@ -110,6 +110,9 @@ class _Kernel(serialize.ReduceMixin):
                if (f'__numba_wrapper_{fn}' in lib.get_asm_str())]
 
         if res:
+            if not config.CUDA_USE_NVIDIA_BINDING:
+                raise NotImplementedError('float16 needs NV Binding')
+
             # Path to the source containing the foreign function
 
             basedir = os.path.dirname(os.path.abspath(__file__))

which causes the example

from numba import cuda, types

@cuda.jit('void()')
def f():
    types.float16(1) / types.float16(2)

to emit:

  File "/home/gmarkall/numbadev/numba/numba/cuda/dispatcher.py", line 114, in __init__
    raise NotImplementedError('float16 needs NV Binding')
NotImplementedError: float16 needs NV Binding

?

testhound · 2023-02-15T17:39:06Z

@gmarkall I was thinking about that after I mentioned "higher up in the stack", and yes that works as well. I'll make that change.

testhound · 2023-02-16T06:09:47Z

gpuci run tests

testhound · 2023-02-16T06:24:40Z

@gmarkall error message updated.

gmarkall · 2023-02-16T10:22:33Z

    def test_linking_cu_ctypes_unsupported(self):
-        msg = ('Linking CUDA source files is not supported with the ctypes '
-               'binding')
+        msg = ('Support for fp16 or linking CUDA source files '


The expected message doesn't seem to match the message added in dispatcher.py. Shouldn't this test be failing?

@gmarkall the test was not failing it appears because it was not being run due to the skip_with_cuda_python decorator.

gmarkall · 2023-02-16T10:22:57Z

@testhound Some questions on the diff.

…config variable for proper testing

testhound · 2023-02-16T21:18:13Z

gpuci run tests

testhound · 2023-02-16T22:31:45Z

gpuci run tests

testhound · 2023-02-17T03:56:22Z

@gmarkall test case updates made to address your comments.

gmarkall

Thanks for the changes. Whilst running test_linking_cu_ctypes_unsupported in a subprocess is a positive change (so it's good to keep it in), my point was that it doesn't test the changes you're making in this PR, which raises when the automatic linking of a .cufile is induced by the use of a float16 function that requires it - so, another test needs to be added for the new exception that's added here.

There are some other suggestions on the diff relating to spelling errors and simplification.

gmarkall · 2023-02-17T11:55:33Z

+                msg = ("Use of float16 requires the use of the NVIDIA CUDA "
+                       "bindings and settng the "
+                       "NUMBA_CUDA_USE_NVIDIA_BINDING environment variable to "
+                       "1. Relevant documentation is available here:\n"
+                       f"{s}")


Some simplification of this comment (and the typo fix suggested yesterday incorporated):

Suggested change

msg = ("Use of float16 requires the use of the NVIDIA CUDA "

"bindings and settng the "

"NUMBA_CUDA_USE_NVIDIA_BINDING environment variable to "

"1. Relevant documentation is available here:\n"

f"{s}")

msg = ("Use of float16 requires the NVIDIA CUDA bindings and "

"setting the NUMBA_CUDA_USE_NVIDIA_BINDING environment "

f"variable to 1. See {s}")

gmarkall · 2023-02-17T11:56:04Z

        self.assertIn('in the compilation of "error.cu"', msg)

-    @skip_with_cuda_python
+    # We need to run the test ina subprocess because the Linker class


Suggested change

# We need to run the test ina subprocess because the Linker class

# We need to run the test in a subprocess because the Linker class

gmarkall · 2023-03-17T17:47:50Z

@stuartarchibald Could this have a CUDA BF run please?

gmarkall

This looks good - just needs a BF run to move to RTM.

sklam · 2023-03-27T22:23:32Z

BFID numba_smoketest_cuda_yaml_192

sklam · 2023-03-27T22:37:47Z

Buildfarm failed with missing jitlink.cu

======================================================================
FAIL: test_linking_cu_ctypes_unsupported (numba.cuda.tests.cudadrv.test_linker.TestLinker)
----------------------------------------------------------------------
Traceback (most recent call last):
File "/opt/conda/envs/testenv_3e4abae3/lib/python3.9/site-packages/numba/tests/support.py", line 644, in inner
self.subprocess_test_runner(test_module=self.__module__,
File "/opt/conda/envs/testenv_3e4abae3/lib/python3.9/site-packages/numba/tests/support.py", line 629, in subprocess_test_runner
self.assertEqual(status.returncode, 0, streams)
AssertionError: 1 != 0 :
captured stdout:
captured stderr: E
======================================================================
ERROR: test_linking_cu_ctypes_unsupported (numba.cuda.tests.cudadrv.test_linker.TestLinker)
----------------------------------------------------------------------
Traceback (most recent call last):
File "/opt/conda/envs/testenv_3e4abae3/lib/python3.9/site-packages/numba/tests/support.py", line 652, in inner
func(self)
File "/opt/conda/envs/testenv_3e4abae3/lib/python3.9/site-packages/numba/cuda/tests/cudadrv/test_linker.py", line 223, in test_linking_cu_ctypes_unsupported
def f():
File "/opt/conda/envs/testenv_3e4abae3/lib/python3.9/site-packages/numba/cuda/decorators.py", line 133, in _jit
disp.compile(argtypes)
File "/opt/conda/envs/testenv_3e4abae3/lib/python3.9/site-packages/numba/cuda/dispatcher.py", line 929, in compile
kernel.bind()
File "/opt/conda/envs/testenv_3e4abae3/lib/python3.9/site-packages/numba/cuda/dispatcher.py", line 208, in bind
self._codelibrary.get_cufunc()
File "/opt/conda/envs/testenv_3e4abae3/lib/python3.9/site-packages/numba/cuda/codegen.py", line 179, in get_cufunc
cubin = self.get_cubin(cc=device.compute_capability)
File "/opt/conda/envs/testenv_3e4abae3/lib/python3.9/site-packages/numba/cuda/codegen.py", line 158, in get_cubin
linker.add_file_guess_ext(path)
File "/opt/conda/envs/testenv_3e4abae3/lib/python3.9/site-packages/numba/cuda/cudadrv/driver.py", line 2652, in add_file_guess_ext
self.add_cu_file(path)
File "/opt/conda/envs/testenv_3e4abae3/lib/python3.9/site-packages/numba/cuda/cudadrv/driver.py", line 2642, in add_cu_file
with open(path, 'rb') as f:
FileNotFoundError: [Errno 2] No such file or directory: 'numba/cuda/tests/data/jitlink.cu'

testhound · 2023-03-27T23:00:54Z

Buildfarm failed with missing jitlink.cu

======================================================================
FAIL: test_linking_cu_ctypes_unsupported (numba.cuda.tests.cudadrv.test_linker.TestLinker)
----------------------------------------------------------------------
Traceback (most recent call last):
File "/opt/conda/envs/testenv_3e4abae3/lib/python3.9/site-packages/numba/tests/support.py", line 644, in inner
self.subprocess_test_runner(test_module=self.__module__,
File "/opt/conda/envs/testenv_3e4abae3/lib/python3.9/site-packages/numba/tests/support.py", line 629, in subprocess_test_runner
self.assertEqual(status.returncode, 0, streams)
AssertionError: 1 != 0 :
captured stdout:
captured stderr: E
======================================================================
ERROR: test_linking_cu_ctypes_unsupported (numba.cuda.tests.cudadrv.test_linker.TestLinker)
----------------------------------------------------------------------
Traceback (most recent call last):
File "/opt/conda/envs/testenv_3e4abae3/lib/python3.9/site-packages/numba/tests/support.py", line 652, in inner
func(self)
File "/opt/conda/envs/testenv_3e4abae3/lib/python3.9/site-packages/numba/cuda/tests/cudadrv/test_linker.py", line 223, in test_linking_cu_ctypes_unsupported
def f():
File "/opt/conda/envs/testenv_3e4abae3/lib/python3.9/site-packages/numba/cuda/decorators.py", line 133, in _jit
disp.compile(argtypes)
File "/opt/conda/envs/testenv_3e4abae3/lib/python3.9/site-packages/numba/cuda/dispatcher.py", line 929, in compile
kernel.bind()
File "/opt/conda/envs/testenv_3e4abae3/lib/python3.9/site-packages/numba/cuda/dispatcher.py", line 208, in bind
self._codelibrary.get_cufunc()
File "/opt/conda/envs/testenv_3e4abae3/lib/python3.9/site-packages/numba/cuda/codegen.py", line 179, in get_cufunc
cubin = self.get_cubin(cc=device.compute_capability)
File "/opt/conda/envs/testenv_3e4abae3/lib/python3.9/site-packages/numba/cuda/codegen.py", line 158, in get_cubin
linker.add_file_guess_ext(path)
File "/opt/conda/envs/testenv_3e4abae3/lib/python3.9/site-packages/numba/cuda/cudadrv/driver.py", line 2652, in add_file_guess_ext
self.add_cu_file(path)
File "/opt/conda/envs/testenv_3e4abae3/lib/python3.9/site-packages/numba/cuda/cudadrv/driver.py", line 2642, in add_cu_file
with open(path, 'rb') as f:
FileNotFoundError: [Errno 2] No such file or directory: 'numba/cuda/tests/data/jitlink.cu'

@sklam this error confuses me as I explicitly fixed the test case to refer to the new location of jitlink.cu and this passes all CI tests which refers to this location.

…rr_msg

gmarkall · 2023-03-28T08:56:42Z

gpuci run tests

gmarkall · 2023-03-28T08:57:52Z

@sklam this error confuses me as I explicitly fixed the test case to refer to the new location of jitlink.cu and this passes all CI tests which refers to this location.

The test contained a hardcoded expectation of the working directory, which doesn't hold in all cases (in the buildfarm in this instance). Test data locations should be determined using test_data_dir. As we're close to the wire for 0.57 I've pushed a fix for this branch.

gmarkall · 2023-03-28T09:48:59Z

I've tested this locally on both Linux and Windows, both in the Numba tree (python runtests.py ...) and outside of it (python -m numba.runtests ...) so I'm reasonably confident this should be OK on the buildfarm now. Could this have another buildfarm run please?

sklam · 2023-03-28T13:55:51Z

new BFID: numba_smoketest_cuda_yaml_193

sklam · 2023-03-28T15:28:03Z

new BFID: numba_smoketest_cuda_yaml_193

passed

testhound · 2023-03-28T15:31:27Z

@sklam this error confuses me as I explicitly fixed the test case to refer to the new location of jitlink.cu and this passes all CI tests which refers to this location.

The test contained a hardcoded expectation of the working directory, which doesn't hold in all cases (in the buildfarm in this instance). Test data locations should be determined using test_data_dir. As we're close to the wire for 0.57 I've pushed a fix for this branch.

@gmarkall thank you for the fix.

gmarkall

I think this is good to merge.

gmarkall · 2023-03-28T15:54:35Z

gpuci run tests

gmarkall · 2023-03-28T15:55:19Z

Not sure what happened to gpuci, just running again to make sure there's nothing odd lurking.

gmarkall · 2023-03-28T16:20:34Z

Marking as RTM as gpuCI had no surprises.

Update error message to include f16 usage

aec2c5d

testhound requested a review from gmarkall as a code owner February 8, 2023 07:10

Fix flake8 error

12288b8

Merge branch 'main' into testhound/cu_linking_err_msg

369b016

gmarkall requested changes Feb 14, 2023

View reviewed changes

gmarkall mentioned this pull request Feb 14, 2023

Add new method is_fp16_supported #8634

Merged

Update error message to be more informative

7704080

gmarkall added the 4 - Waiting on author Waiting for author to respond to review label Feb 15, 2023

Separate out float16 error exception from cu linking

db210b6

gmarkall reviewed Feb 16, 2023

View reviewed changes

Comment thread numba/cuda/dispatcher.py Outdated

gmarkall reviewed Feb 16, 2023

View reviewed changes

gmarkall added the CUDA CUDA related issue/PR label Feb 16, 2023

Run test case unconditionally, fix file not being found and override …

275243e

…config variable for proper testing

gmarkall requested changes Feb 17, 2023

View reviewed changes

stuartarchibald assigned gmarkall Feb 24, 2023

gmarkall added this to the Numba 0.57 RC milestone Mar 17, 2023

gmarkall added 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 Mar 17, 2023

gmarkall previously approved these changes Mar 27, 2023

View reviewed changes

sklam removed the 4 - Waiting on CI Review etc done, waiting for CI to finish label Mar 27, 2023

gmarkall added 2 commits March 28, 2023 09:39

Merge remote-tracking branch 'numba/main' into testhound/cu_linking_e…

51f914d

…rr_msg

CUDA test_linker.py: fix jitlink.cu path

18dae59

gmarkall dismissed their stale review via 18dae59 March 28, 2023 08:56

gmarkall added the 4 - Waiting on CI Review etc done, waiting for CI to finish label Mar 28, 2023

gmarkall approved these changes Mar 28, 2023

View reviewed changes

gmarkall added the 5 - Ready to merge Review and testing done, is ready to merge label Mar 28, 2023

sklam merged commit c1276d7 into numba:main Mar 28, 2023

	# We need to run the test ina subprocess because the Linker class
	# We need to run the test in a subprocess because the Linker class

Conversation

testhound commented Feb 8, 2023

Uh oh!

testhound commented Feb 8, 2023

Uh oh!

testhound commented Feb 13, 2023

Uh oh!

gmarkall commented Feb 13, 2023

Uh oh!

gmarkall left a comment

Choose a reason for hiding this comment

Uh oh!

gmarkall Feb 14, 2023

Choose a reason for hiding this comment

Uh oh!

testhound commented Feb 15, 2023

Uh oh!

testhound commented Feb 15, 2023

Uh oh!

gmarkall commented Feb 15, 2023

Uh oh!

testhound commented Feb 15, 2023

Uh oh!

gmarkall commented Feb 15, 2023

Uh oh!

testhound commented Feb 15, 2023

Uh oh!

testhound commented Feb 16, 2023

Uh oh!

testhound commented Feb 16, 2023

Uh oh!

Uh oh!

gmarkall Feb 16, 2023

Choose a reason for hiding this comment

Uh oh!

testhound Feb 16, 2023

Choose a reason for hiding this comment

Uh oh!

gmarkall commented Feb 16, 2023

Uh oh!

testhound commented Feb 16, 2023

Uh oh!

testhound commented Feb 16, 2023

Uh oh!

testhound commented Feb 17, 2023

Uh oh!

gmarkall left a comment

Choose a reason for hiding this comment

Uh oh!

gmarkall Feb 17, 2023

Choose a reason for hiding this comment

Uh oh!

gmarkall Feb 17, 2023

Choose a reason for hiding this comment

Uh oh!

gmarkall commented Mar 17, 2023

Uh oh!

gmarkall left a comment

Choose a reason for hiding this comment

Uh oh!

sklam commented Mar 27, 2023

Uh oh!

sklam commented Mar 27, 2023

Uh oh!

testhound commented Mar 27, 2023

Uh oh!

gmarkall commented Mar 28, 2023

Uh oh!

gmarkall commented Mar 28, 2023

Uh oh!

gmarkall commented Mar 28, 2023

Uh oh!

sklam commented Mar 28, 2023

Uh oh!

sklam commented Mar 28, 2023

Uh oh!

testhound commented Mar 28, 2023

Uh oh!

gmarkall left a comment

Choose a reason for hiding this comment