CUDA: Remove items deprecated in 0.53 + simulator test fixes

[gmarkall]

This PR removes items deprecated in 0.53 and scheduled for removal in 0.54. These include:

• The compute_capabilty kwarg to inspect methods.
• The ability to index dicts returns by inspect method using a tuple of compute capability and argument types.
• The return of a single item (instead of a dict) from the inspection methods and ptx property of specialized dispatchers.
• The definition and definitions property of dispatchers.

Whilst working on this I discovered a couple of other issues in the simulator:

• The testing of cooperative groups was accidentally skipped because skip_if_cudadevrt_missing always skipped tests on the simulator - this is now resolved.
• Some tests inspected some fake PTX from the simulator - this makes no sense, so the ptx property of fake kernels has been removed, and the tests which inspected it are now skipped on cudasim.

Some other remarks:

• Various tests made use of the deprecated functionality - these have now been updated to use the current interfaces.
• Relevant documentation (the Cooperative groups example and the deprecation notices) is updated.
• Tests of the deprecation notices have been removed accordingly.

[stuartarchibald]

Thanks for the patch, great to see this cleaned up. There's a few minor things to look at and once resolved this will be good to merge. Thanks again!

[gmarkall]

@stuartarchibald Many thanks for the review - comments should be addressed now.

[gmarkall]

Is the fail a Frosty Thompson?

INFO: Running slice of discovered tests: (7,None,22)
cc1plus: warning: command line option '-Wstrict-prototypes' is valid for C/ObjC but not for C++
skipped HSA tests
skipped HSA tests
Parallel: 402. Serial: 62
s....................................................................................................................................x...............cc1plus: warning: command line option '-Wstrict-prototypes' is valid for C/ObjC but not for C++
..................................s................ssssssssssssssss..........F...............................................................................................................................................................................sss..ss..........sssssss.s..s....sss..........s............sss
======================================================================
FAIL: test_fork_from_non_main_thread (numba.tests.test_parallel_backend.TestTBBSpecificIssues)
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/vsts/work/1/s/numba/tests/test_parallel_backend.py", line 921, in test_fork_from_non_main_thread
    self.assertIn(msg_head, err)
AssertionError: 'Attempted to fork from a non-main thread, the TBB library' not found in ''

----------------------------------------------------------------------
Ran 464 tests in 159.435s

FAILED (failures=1, skipped=39, expected failures=1)

[stuartarchibald]

Is the fail a Frosty Thompson?

No, that's the fork from non-main thread problem reappearing, would guess TBB 2021 related. Frosty Thompson appears as a live lock.

[stuartarchibald]

Is the fail a Frosty Thompson?

No, that's the fork from non-main thread problem reappearing, would guess TBB 2021 related. Frosty Thompson appears as a live lock.

Fixed in #6966

[stuartarchibald]

Thanks for the fixes, please could you resolve the conflict against mainline, this can then have a build farm run?

[gmarkall]

/azp run

[azure-pipelines]

Azure Pipelines successfully started running 1 pipeline(s).

[gmarkall]

(Re-running due to conda 500 / 530 errors in the prior test run)

[stuartarchibald]

Thanks for the patch and fixes.

[stuartarchibald]

Buildfarm ID: numba_smoketest_cuda_yaml_61.

[stuartarchibald]

Buildfarm ID: numba_smoketest_cuda_yaml_61.

This failed. All fails are on windows.

• py3.8, np1.17, cuda 9.2
• py3.8, np1.17, cuda 10.0
• py3.9, np1.19, cuda 10.0
• py3.9, np1.19, cuda 9.2

Failure is the same on each:

<snip>
FAIL: test_const_record_optimization (numba.cuda.tests.cudapy.test_constmem.TestCudaConstantMemory)
-----------------------------------------------------------------------------
Trace back:
"numba\cuda\tests\test_constmem.py", line 200, in test_const_record_optimization
    self.assertGreaterEqual(u8_load_count, expected_load_count,
AssertionError: 11 not greater than or equal to 12: load record values as individual bytes
<snip>

[gmarkall]

@stuartarchibald Thanks for the report. I've removed the failing test and a similar one, for the rationale outlined in the commit message - in summary I think these tests are a bit unhelpful / nonsensical and have adjusted them several times over the course of various PRs. Explanation from the commit:

These tests check for very specific patterns in the generated PTX that
are unstable across toolkit versions, platforms, and other minor
changes. Since these frequently need adjustment to match the generated
PTX, and have never revealed an underlying issue, is seems more sensible
to remove them instead of continually fixing them up in the presence of
variations.

Note that the latest issue was detected in the PR numba#6840 buildfarm run,
with failures in test_const_record_optimization on Windows for CUDA 9.2
and 10.0. These failures arose because a change to compile the kernel
with an explicit signature instead of specializing the kernel for given
arguments changed the types the kernel was being compiled for on Windows
- the default int on Windows is different to the default int on Linux.
So this test wasn't even testing the same thing on Windows as on Linux,
and this went unnoticed for a very long time!

[stuartarchibald]

@stuartarchibald Thanks for the report. I've removed the failing test and a similar one, for the rationale outlined in the commit message - in summary I think these tests are a bit unhelpful / nonsensical and have adjusted them several times over the course of various PRs. Explanation from the commit:

These tests check for very specific patterns in the generated PTX that
are unstable across toolkit versions, platforms, and other minor
changes. Since these frequently need adjustment to match the generated
PTX, and have never revealed an underlying issue, is seems more sensible
to remove them instead of continually fixing them up in the presence of
variations.

Note that the latest issue was detected in the PR numba#6840 buildfarm run,
with failures in test_const_record_optimization on Windows for CUDA 9.2
and 10.0. These failures arose because a change to compile the kernel
with an explicit signature instead of specializing the kernel for given
arguments changed the types the kernel was being compiled for on Windows
- the default int on Windows is different to the default int on Linux.
So this test wasn't even testing the same thing on Windows as on Linux,
and this went unnoticed for a very long time!

Thanks for the explanation @gmarkall. I agree, if these tests are not really adding anything but very susceptible to external/environmental changes then removing them seems to be the best solution.

[stuartarchibald]

Thanks for the fix.

[stuartarchibald]

Buildfarm ID: numba_smoketest_cuda_yaml_63.

[stuartarchibald]

Buildfarm ID: numba_smoketest_cuda_yaml_63.

This failed due to the issue in #7052. Once that's merged will re-run the build.

[gmarkall]

@stuartarchibald / @esc Could this have another run now please?

[stuartarchibald]

@stuartarchibald / @esc Could this have another run now please?

Buildfarm ID: numba_smoketest_cuda_yaml_67.

[stuartarchibald]

@stuartarchibald / @esc Could this have another run now please?

Buildfarm ID: numba_smoketest_cuda_yaml_67.

Passed.