Introduce cuda.cooperative overloads not requiring temporary storage (NVIDIA#2528)

* Modernize pkg resource query

* Add cooperative overloads without shared memory

* Start fixing temp storage

* Incorporate template params into mangling

* Condence dict access

* Fix temporary storage indexing for sub hw waprs

* Test multiple warps

* Disable alloc API for sub hw warps

Author: gevtushenko

python/cuda_cooperative/cuda/cooperative/experimental/_nvrtc.py

@@ -10,6 +10,7 @@
 import importlib.resources as pkg_resources
 import functools
 
+
 def CHECK_NVRTC(err, prog):
     if err != nvrtc.nvrtcResult.NVRTC_SUCCESS:
         err, logsize = nvrtc.nvrtcGetProgramLogSize(prog)
@@ -39,8 +40,8 @@ def get_cuda_path():
 # rdc is true or false
 # code is lto or ptx
 # @cache
-@functools.lru_cache(maxsize=32) # Always enabled
-@disk_cache # Optional, see caching.py
+@functools.lru_cache(maxsize=32)  # Always enabled
+@disk_cache  # Optional, see caching.py
 def compile_impl(cpp, cc, rdc, code, nvrtc_path, nvrtc_version):
     check_in('rdc', rdc, [True, False])
     check_in('code', code, ['lto', 'ptx'])
@@ -54,11 +55,11 @@ def compile_impl(cpp, cc, rdc, code, nvrtc_path, nvrtc_version):
         libcudacxx_path = os.path.join(include_path, 'libcudacxx')
         cuda_include_path = os.path.join(get_cuda_path(), 'include')
 
-    opts = [b"--std=c++17", \
-            bytes(f"--include-path={cub_path}", encoding='ascii'), \
-            bytes(f"--include-path={thrust_path}", encoding='ascii'), \
-            bytes(f"--include-path={libcudacxx_path}", encoding='ascii'), \
-            bytes(f"--include-path={cuda_include_path}", encoding='ascii'), \
+    opts = [b"--std=c++17",
+            bytes(f"--include-path={cub_path}", encoding='ascii'),
+            bytes(f"--include-path={thrust_path}", encoding='ascii'),
+            bytes(f"--include-path={libcudacxx_path}", encoding='ascii'),
+            bytes(f"--include-path={cuda_include_path}", encoding='ascii'),
             bytes(f"--gpu-architecture=compute_{cc}", encoding='ascii')]
     if rdc:
         opts += [b"--relocatable-device-code=true"]
@@ -70,7 +71,8 @@ def compile_impl(cpp, cc, rdc, code, nvrtc_path, nvrtc_version):
     opts += [b"-DCCCL_DISABLE_BF16_SUPPORT"]
 
     # Create program
-    err, prog = nvrtc.nvrtcCreateProgram(str.encode(cpp), b"code.cu", 0, [], [])
+    err, prog = nvrtc.nvrtcCreateProgram(
+        str.encode(cpp), b"code.cu", 0, [], [])
     if err != nvrtc.nvrtcResult.NVRTC_SUCCESS:
         raise RuntimeError(f"nvrtcCreateProgram error: {err}")
 
@@ -103,12 +105,13 @@ def compile_impl(cpp, cc, rdc, code, nvrtc_path, nvrtc_version):
 
         return ptx.decode('ascii')
 
+
 def compile(**kwargs):
 
     err, major, minor = nvrtc.nvrtcVersion()
     if err != nvrtc.nvrtcResult.NVRTC_SUCCESS:
         raise RuntimeError(f"nvrtcVersion error: {err}")
     nvrtc_version = version(major, minor)
-    return nvrtc_version, compile_impl(**kwargs, \
-                        nvrtc_path=nvrtc.__file__, \
-                        nvrtc_version=nvrtc_version)
+    return nvrtc_version, compile_impl(**kwargs,
+                                       nvrtc_path=nvrtc.__file__,
+                                       nvrtc_version=nvrtc_version)