diff --git a/cuda_core/cuda/core/utils/_program_cache/_keys.py b/cuda_core/cuda/core/utils/_program_cache/_keys.py index 2df2d893835..2919f48ee34 100644 --- a/cuda_core/cuda/core/utils/_program_cache/_keys.py +++ b/cuda_core/cuda/core/utils/_program_cache/_keys.py @@ -611,12 +611,13 @@ def hash_extra_payload(self, options: ProgramOptions, update: Callable[[str, byt if isinstance(item, (tuple, list)) and len(item) == 2: name, src = item update("extra_source_name", str(name).encode("utf-8")) - if isinstance(src, str): - update("extra_source_code", src.encode("utf-8")) - elif isinstance(src, (bytes, bytearray)): - update("extra_source_code", bytes(src)) - else: - update("extra_source_code", str(src).encode("utf-8")) + match src: + case str(): + update("extra_source_code", src.encode("utf-8")) + case bytes() | bytearray(): + update("extra_source_code", bytes(src)) + case _: + update("extra_source_code", str(src).encode("utf-8")) else: # Fallback for unexpected format. update("extra_source", str(item).encode("utf-8")) diff --git a/samples/cuda_bindings/3_CUDA_Features/globalToShmemAsyncCopy/globalToShmemAsyncCopy.py b/samples/cuda_bindings/3_CUDA_Features/globalToShmemAsyncCopy/globalToShmemAsyncCopy.py index f716699aab9..7d4b436fe1c 100644 --- a/samples/cuda_bindings/3_CUDA_Features/globalToShmemAsyncCopy/globalToShmemAsyncCopy.py +++ b/samples/cuda_bindings/3_CUDA_Features/globalToShmemAsyncCopy/globalToShmemAsyncCopy.py @@ -857,134 +857,135 @@ def matrix_multiply(dims_a, dims_b, kernel_number): (d_c, d_a, d_b, dims_a.x, dims_b.x), (ctypes.c_void_p, ctypes.c_void_p, ctypes.c_void_p, ctypes.c_int, ctypes.c_int), ) - if kernel_number == Kernels.AsyncCopyMultiStageLargeChunk: - check_cuda_errors( - cuda.cuLaunchKernel( - _MatrixMulAsyncCopyMultiStageLargeChunk, - grid.x, - grid.y, - grid.z, # grid dim - threads.x, - threads.y, - threads.z, # block dim - 0, # shared mem - stream, # stream - kernel_arguments, - 0, - ) - ) # arguments - elif kernel_number == Kernels.AsyncCopyLargeChunk: - check_cuda_errors( - cuda.cuLaunchKernel( - _MatrixMulAsyncCopyLargeChunk, - grid.x, - grid.y, - grid.z, # grid dim - threads.x, - threads.y, - threads.z, # block dim - 0, # shared mem - stream, # stream - kernel_arguments, - 0, - ) - ) # arguments - elif kernel_number == Kernels.AsyncCopyLargeChunkAWBarrier: - check_cuda_errors( - cuda.cuLaunchKernel( - _MatrixMulAsyncCopyLargeChunkAWBarrier, - grid.x, - grid.y, - grid.z, # grid dim - threads.x, - threads.y, - threads.z, # block dim - 0, # shared mem - stream, # stream - kernel_arguments, - 0, - ) - ) # arguments - elif kernel_number == Kernels.AsyncCopyMultiStageSharedState: - check_cuda_errors( - cuda.cuLaunchKernel( - _MatrixMulAsyncCopyMultiStageSharedState, - grid_shared_state_kernel.x, - grid_shared_state_kernel.y, - grid_shared_state_kernel.z, # grid dim - threads_shared_state_kernel.x, - threads_shared_state_kernel.y, - threads_shared_state_kernel.z, # block dim - 0, # shared mem - stream, # stream - kernel_arguments, - 0, - ) - ) # arguments - elif kernel_number == Kernels.AsyncCopyMultiStage: - check_cuda_errors( - cuda.cuLaunchKernel( - _MatrixMulAsyncCopyMultiStage, - grid.x, - grid.y, - grid.z, # grid dim - threads.x, - threads.y, - threads.z, # block dim - 0, # shared mem - stream, # stream - kernel_arguments, - 0, - ) - ) # arguments - elif kernel_number == Kernels.AsyncCopySingleStage: - check_cuda_errors( - cuda.cuLaunchKernel( - _MatrixMulAsyncCopySingleStage, - grid.x, - grid.y, - grid.z, # grid dim - threads.x, - threads.y, - threads.z, # block dim - 0, # shared mem - stream, # stream - kernel_arguments, - 0, - ) - ) # arguments - elif kernel_number == Kernels.Naive: - check_cuda_errors( - cuda.cuLaunchKernel( - _MatrixMulNaive, - grid.x, - grid.y, - grid.z, # grid dim - threads.x, - threads.y, - threads.z, # block dim - 0, # shared mem - stream, # stream - kernel_arguments, - 0, - ) - ) # arguments - elif kernel_number == Kernels.NaiveLargeChunk: - check_cuda_errors( - cuda.cuLaunchKernel( - _MatrixMulNaiveLargeChunk, - grid.x, - grid.y, - grid.z, # grid dim - threads.x, - threads.y, - threads.z, # block dim - 0, # shared mem - stream, # stream - kernel_arguments, - 0, - ) - ) # arguments + match kernel_number: + case Kernels.AsyncCopyMultiStageLargeChunk: + check_cuda_errors( + cuda.cuLaunchKernel( + _MatrixMulAsyncCopyMultiStageLargeChunk, + grid.x, + grid.y, + grid.z, # grid dim + threads.x, + threads.y, + threads.z, # block dim + 0, # shared mem + stream, # stream + kernel_arguments, + 0, + ) + ) # arguments + case Kernels.AsyncCopyLargeChunk: + check_cuda_errors( + cuda.cuLaunchKernel( + _MatrixMulAsyncCopyLargeChunk, + grid.x, + grid.y, + grid.z, # grid dim + threads.x, + threads.y, + threads.z, # block dim + 0, # shared mem + stream, # stream + kernel_arguments, + 0, + ) + ) # arguments + case Kernels.AsyncCopyLargeChunkAWBarrier: + check_cuda_errors( + cuda.cuLaunchKernel( + _MatrixMulAsyncCopyLargeChunkAWBarrier, + grid.x, + grid.y, + grid.z, # grid dim + threads.x, + threads.y, + threads.z, # block dim + 0, # shared mem + stream, # stream + kernel_arguments, + 0, + ) + ) # arguments + case Kernels.AsyncCopyMultiStageSharedState: + check_cuda_errors( + cuda.cuLaunchKernel( + _MatrixMulAsyncCopyMultiStageSharedState, + grid_shared_state_kernel.x, + grid_shared_state_kernel.y, + grid_shared_state_kernel.z, # grid dim + threads_shared_state_kernel.x, + threads_shared_state_kernel.y, + threads_shared_state_kernel.z, # block dim + 0, # shared mem + stream, # stream + kernel_arguments, + 0, + ) + ) # arguments + case Kernels.AsyncCopyMultiStage: + check_cuda_errors( + cuda.cuLaunchKernel( + _MatrixMulAsyncCopyMultiStage, + grid.x, + grid.y, + grid.z, # grid dim + threads.x, + threads.y, + threads.z, # block dim + 0, # shared mem + stream, # stream + kernel_arguments, + 0, + ) + ) # arguments + case Kernels.AsyncCopySingleStage: + check_cuda_errors( + cuda.cuLaunchKernel( + _MatrixMulAsyncCopySingleStage, + grid.x, + grid.y, + grid.z, # grid dim + threads.x, + threads.y, + threads.z, # block dim + 0, # shared mem + stream, # stream + kernel_arguments, + 0, + ) + ) # arguments + case Kernels.Naive: + check_cuda_errors( + cuda.cuLaunchKernel( + _MatrixMulNaive, + grid.x, + grid.y, + grid.z, # grid dim + threads.x, + threads.y, + threads.z, # block dim + 0, # shared mem + stream, # stream + kernel_arguments, + 0, + ) + ) # arguments + case Kernels.NaiveLargeChunk: + check_cuda_errors( + cuda.cuLaunchKernel( + _MatrixMulNaiveLargeChunk, + grid.x, + grid.y, + grid.z, # grid dim + threads.x, + threads.y, + threads.z, # block dim + 0, # shared mem + stream, # stream + kernel_arguments, + 0, + ) + ) # arguments check_cuda_errors(cudart.cudaStreamSynchronize(stream)) @@ -994,134 +995,135 @@ def matrix_multiply(dims_a, dims_b, kernel_number): # Record the start event check_cuda_errors(cudart.cudaEventRecord(start, stream)) - if kernel_number == Kernels.AsyncCopyMultiStageLargeChunk: - check_cuda_errors( - cuda.cuLaunchKernel( - _MatrixMulAsyncCopyMultiStageLargeChunk, - grid.x, - grid.y, - grid.z, # grid dim - threads.x, - threads.y, - threads.z, # block dim - 0, # shared mem - stream, # stream - kernel_arguments, - 0, - ) - ) # arguments - elif kernel_number == Kernels.AsyncCopyLargeChunk: - check_cuda_errors( - cuda.cuLaunchKernel( - _MatrixMulAsyncCopyLargeChunk, - grid.x, - grid.y, - grid.z, # grid dim - threads.x, - threads.y, - threads.z, # block dim - 0, # shared mem - stream, # stream - kernel_arguments, - 0, - ) - ) # arguments - elif kernel_number == Kernels.AsyncCopyLargeChunkAWBarrier: - check_cuda_errors( - cuda.cuLaunchKernel( - _MatrixMulAsyncCopyLargeChunkAWBarrier, - grid.x, - grid.y, - grid.z, # grid dim - threads.x, - threads.y, - threads.z, # block dim - 0, # shared mem - stream, # stream - kernel_arguments, - 0, - ) - ) # arguments - elif kernel_number == Kernels.AsyncCopyMultiStageSharedState: - check_cuda_errors( - cuda.cuLaunchKernel( - _MatrixMulAsyncCopyMultiStageSharedState, - grid_shared_state_kernel.x, - grid_shared_state_kernel.y, - grid_shared_state_kernel.z, # grid dim - threads_shared_state_kernel.x, - threads_shared_state_kernel.y, - threads_shared_state_kernel.z, # block dim - 0, # shared mem - stream, # stream - kernel_arguments, - 0, - ) - ) # arguments - elif kernel_number == Kernels.AsyncCopyMultiStage: - check_cuda_errors( - cuda.cuLaunchKernel( - _MatrixMulAsyncCopyMultiStage, - grid.x, - grid.y, - grid.z, # grid dim - threads.x, - threads.y, - threads.z, # block dim - 0, # shared mem - stream, # stream - kernel_arguments, - 0, - ) - ) # arguments - elif kernel_number == Kernels.AsyncCopySingleStage: - check_cuda_errors( - cuda.cuLaunchKernel( - _MatrixMulAsyncCopySingleStage, - grid.x, - grid.y, - grid.z, # grid dim - threads.x, - threads.y, - threads.z, # block dim - 0, # shared mem - stream, # stream - kernel_arguments, - 0, - ) - ) # arguments - elif kernel_number == Kernels.Naive: - check_cuda_errors( - cuda.cuLaunchKernel( - _MatrixMulNaive, - grid.x, - grid.y, - grid.z, # grid dim - threads.x, - threads.y, - threads.z, # block dim - 0, # shared mem - stream, # stream - kernel_arguments, - 0, - ) - ) # arguments - elif kernel_number == Kernels.NaiveLargeChunk: - check_cuda_errors( - cuda.cuLaunchKernel( - _MatrixMulNaiveLargeChunk, - grid.x, - grid.y, - grid.z, # grid dim - threads.x, - threads.y, - threads.z, # block dim - 0, # shared mem - stream, # stream - kernel_arguments, - 0, - ) - ) # arguments + match kernel_number: + case Kernels.AsyncCopyMultiStageLargeChunk: + check_cuda_errors( + cuda.cuLaunchKernel( + _MatrixMulAsyncCopyMultiStageLargeChunk, + grid.x, + grid.y, + grid.z, # grid dim + threads.x, + threads.y, + threads.z, # block dim + 0, # shared mem + stream, # stream + kernel_arguments, + 0, + ) + ) # arguments + case Kernels.AsyncCopyLargeChunk: + check_cuda_errors( + cuda.cuLaunchKernel( + _MatrixMulAsyncCopyLargeChunk, + grid.x, + grid.y, + grid.z, # grid dim + threads.x, + threads.y, + threads.z, # block dim + 0, # shared mem + stream, # stream + kernel_arguments, + 0, + ) + ) # arguments + case Kernels.AsyncCopyLargeChunkAWBarrier: + check_cuda_errors( + cuda.cuLaunchKernel( + _MatrixMulAsyncCopyLargeChunkAWBarrier, + grid.x, + grid.y, + grid.z, # grid dim + threads.x, + threads.y, + threads.z, # block dim + 0, # shared mem + stream, # stream + kernel_arguments, + 0, + ) + ) # arguments + case Kernels.AsyncCopyMultiStageSharedState: + check_cuda_errors( + cuda.cuLaunchKernel( + _MatrixMulAsyncCopyMultiStageSharedState, + grid_shared_state_kernel.x, + grid_shared_state_kernel.y, + grid_shared_state_kernel.z, # grid dim + threads_shared_state_kernel.x, + threads_shared_state_kernel.y, + threads_shared_state_kernel.z, # block dim + 0, # shared mem + stream, # stream + kernel_arguments, + 0, + ) + ) # arguments + case Kernels.AsyncCopyMultiStage: + check_cuda_errors( + cuda.cuLaunchKernel( + _MatrixMulAsyncCopyMultiStage, + grid.x, + grid.y, + grid.z, # grid dim + threads.x, + threads.y, + threads.z, # block dim + 0, # shared mem + stream, # stream + kernel_arguments, + 0, + ) + ) # arguments + case Kernels.AsyncCopySingleStage: + check_cuda_errors( + cuda.cuLaunchKernel( + _MatrixMulAsyncCopySingleStage, + grid.x, + grid.y, + grid.z, # grid dim + threads.x, + threads.y, + threads.z, # block dim + 0, # shared mem + stream, # stream + kernel_arguments, + 0, + ) + ) # arguments + case Kernels.Naive: + check_cuda_errors( + cuda.cuLaunchKernel( + _MatrixMulNaive, + grid.x, + grid.y, + grid.z, # grid dim + threads.x, + threads.y, + threads.z, # block dim + 0, # shared mem + stream, # stream + kernel_arguments, + 0, + ) + ) # arguments + case Kernels.NaiveLargeChunk: + check_cuda_errors( + cuda.cuLaunchKernel( + _MatrixMulNaiveLargeChunk, + grid.x, + grid.y, + grid.z, # grid dim + threads.x, + threads.y, + threads.z, # block dim + 0, # shared mem + stream, # stream + kernel_arguments, + 0, + ) + ) # arguments # Record the stop event check_cuda_errors(cudart.cudaEventRecord(stop, stream)) diff --git a/samples/cuda_bindings/4_CUDA_Libraries/nvidiaSmi/nvidiaSmi.py b/samples/cuda_bindings/4_CUDA_Libraries/nvidiaSmi/nvidiaSmi.py index 6287138cc49..e6dbb4726b3 100644 --- a/samples/cuda_bindings/4_CUDA_Libraries/nvidiaSmi/nvidiaSmi.py +++ b/samples/cuda_bindings/4_CUDA_Libraries/nvidiaSmi/nvidiaSmi.py @@ -241,14 +241,15 @@ def collect_info(): except nvml.NvmlError: comp_str = "N/A" else: - if compute_mode == nvml.ComputeMode.COMPUTEMODE_DEFAULT: - comp_str = "Default" - elif compute_mode == nvml.ComputeMode.COMPUTEMODE_EXCLUSIVE_PROCESS: - comp_str = "E. Process" - elif compute_mode == nvml.ComputeMode.COMPUTEMODE_PROHIBITED: - comp_str = "Prohibited" - else: - comp_str = "Unknown" + match compute_mode: + case nvml.ComputeMode.COMPUTEMODE_DEFAULT: + comp_str = "Default" + case nvml.ComputeMode.COMPUTEMODE_EXCLUSIVE_PROCESS: + comp_str = "E. Process" + case nvml.ComputeMode.COMPUTEMODE_PROHIBITED: + comp_str = "Prohibited" + case _: + comp_str = "Unknown" device["compute_mode"] = comp_str devices.append(device) diff --git a/samples/cuda_core/glInteropMipmapLod/glInteropMipmapLod.py b/samples/cuda_core/glInteropMipmapLod/glInteropMipmapLod.py index 93808aacbb6..6d4b4da40c6 100644 --- a/samples/cuda_core/glInteropMipmapLod/glInteropMipmapLod.py +++ b/samples/cuda_core/glInteropMipmapLod/glInteropMipmapLod.py @@ -576,13 +576,14 @@ def on_mouse_scroll(_x, _y, _scroll_x, scroll_y): @window.event def on_key_press(symbol, _modifiers): key = pyglet.window.key - if symbol == key.BRACKETLEFT: - state["lod_bias"] = max(-float(num_levels), state["lod_bias"] - LOD_BIAS_STEP) - elif symbol == key.BRACKETRIGHT: - state["lod_bias"] = min(float(num_levels), state["lod_bias"] + LOD_BIAS_STEP) - elif symbol == key.R: - state["zoom"] = 1.0 - state["lod_bias"] = 0.0 + match symbol: + case key.BRACKETLEFT: + state["lod_bias"] = max(-float(num_levels), state["lod_bias"] - LOD_BIAS_STEP) + case key.BRACKETRIGHT: + state["lod_bias"] = min(float(num_levels), state["lod_bias"] + LOD_BIAS_STEP) + case key.R: + state["zoom"] = 1.0 + state["lod_bias"] = 0.0 @window.event def on_close(): diff --git a/toolshed/check_cython_abi.py b/toolshed/check_cython_abi.py index b72edf40c5c..93d7cca57e2 100644 --- a/toolshed/check_cython_abi.py +++ b/toolshed/check_cython_abi.py @@ -186,18 +186,19 @@ def _collect_structs_from_tree(node: object) -> dict: result = {} cls = type(node).__name__ - if cls == "CStructOrUnionDefNode": - fields = [] - for attr in node.attributes: - if type(attr).__name__ == "CVarDefNode": - fields.extend(_extract_fields_from_cvardef(attr)) - if fields: - result[node.name] = {"fields": fields} - - elif cls == "CClassDefNode": - fields = _collect_cvardef_fields(node.body) - if fields: - result[node.class_name] = {"fields": fields} + match cls: + case "CStructOrUnionDefNode": + fields = [] + for attr in node.attributes: + if type(attr).__name__ == "CVarDefNode": + fields.extend(_extract_fields_from_cvardef(attr)) + if fields: + result[node.name] = {"fields": fields} + + case "CClassDefNode": + fields = _collect_cvardef_fields(node.body) + if fields: + result[node.class_name] = {"fields": fields} for attr_name in getattr(node, "child_attrs", []): child = getattr(node, attr_name, None)