SASS

Tests for reprospect.testing.binaries.sass.

The tests for the subpackages can be found in:

Extensibility

mypyc compilation can inadvertently seal classes, preventing inheritance. This module verifies that the following types remain subclassable after compilation, allowing users to extend reprospect capabilities:

Warning

This test must be run with the compiled binary distribution of reprospect. Using the source distribution makes tests.testing.binaries.sass.test_extensibility.CannotBeExtended a valid extension.

class tests.testing.binaries.sass.test_extensibility.CannotBeExtendedView on GitHub 

Bases: Fp32AddMatcher

reprospect.testing.binaries.sass.instruction.floating.Fp32AddMatcher was not marked as extensible.

pattern: Final[Pattern[str]]

class tests.testing.binaries.sass.test_extensibility.NewInstructionMatcherView on GitHub 

Bases: InstructionMatcher

Faking a matcher for some unforeseen use case.

match(inst: str | Instruction) → InstructionMatch | NoneView on GitHub 

class tests.testing.binaries.sass.test_extensibility.NewPatternMatcherView on GitHub 

Bases: PatternMatcher

Faking a matcher for some unforeseen use case.

__init__()View on GitHub 

pattern: Final[Pattern[str]]

class tests.testing.binaries.sass.test_extensibility.TestCannotBeExtendedView on GitHub 

Bases: object

test() → NoneView on GitHub 

class tests.testing.binaries.sass.test_extensibility.TestInspectView on GitHub 

Bases: object

Use inspect to retrieve the provenance of objects.

test_Instruction() → NoneView on GitHub 

test_InstructionMatcher() → NoneView on GitHub 

class tests.testing.binaries.sass.test_extensibility.TestInstructionMatchingView on GitHub 

Bases: object

Match instructions with NewInstructionMatcher and NewPatternMatcher.

CONTROLCODE: Final[ControlCode] = ControlCode(stall_count=5, yield_flag=False, read=7, write=7, wait=[True, False, False, False, False, False], reuse={'A': False, 'B': False, 'C': False, 'D': False})

DADD: Final[Instruction] = Instruction(offset=0, instruction='DADD R4, R4, c[0x0][0x180]', hex='0x0', control=ControlCode(stall_count=5, yield_flag=False, read=7, write=7, wait=[True, False, False, False, False, False], reuse={'A': False, 'B': False, 'C': False, 'D': False}))

DMUL: Final[Instruction] = Instruction(offset=0, instruction='DMUL R6, R6, c[0x0][0x188]', hex='0x1', control=ControlCode(stall_count=5, yield_flag=False, read=7, write=7, wait=[True, False, False, False, False, False], reuse={'A': False, 'B': False, 'C': False, 'D': False}))

NOP: Final[Instruction] = Instruction(offset=0, instruction='NOP', hex='0x2', control=ControlCode(stall_count=5, yield_flag=False, read=7, write=7, wait=[True, False, False, False, False, False], reuse={'A': False, 'B': False, 'C': False, 'D': False}))

test_NewInstructionMatcher() → NoneView on GitHub 

test_NewPatternMatcher() → NoneView on GitHub 

class tests.testing.binaries.sass.test_half.TestFp16AddMatcherView on GitHub 

Bases: object

Tests for reprospect.testing.binaries.sass.instruction.half.Fp16AddMatcher.

CODE_AUTO_PACK: Final[str] = '#include "cuda_fp16.h"\n__global__ void test_packed(__half* __restrict__ const dst, const __half* __restrict__ const src) {\n dst[0] += src[0]; dst[1] += src[1];\n}\n'

CODE_FORCE_PACK: Final[str] = '#include "cuda_fp16.h"\n__global__ void test_packed(__half* __restrict__ const dst, const __half* __restrict__ const src) {\n reinterpret_cast<__half2*>(dst)[0] += reinterpret_cast<const __half2*>(src)[0];\n}\n'

CODE_FP16_TO_FP32: Final[str] = '#include "cuda_fp16.h"\n__global__ void test_fp16_to_fp32(float* __restrict__ const dst, const __half* __restrict__ const src, const unsigned int size)\n{\n const auto index = blockIdx.x * blockDim.x + threadIdx.x;\n if (index < size)\n dst[index] = __half2float(src[index]);\n}\n'

test_any() → NoneView on GitHub 

test_force_pack(request, workdir: Path, parameters: Parameters, cmake_file_api) → NoneView on GitHub : One can force the packing in CODE_FORCE_PACK.

test_fp16_to_fp32(request, workdir: Path, parameters: Parameters, cmake_file_api) → NoneView on GitHub : Use CODE_FP16_TO_FP32 to check that the conversion of __half to float always goes through a HADD2.F32 instruction.

test_individual() → NoneView on GitHub 

test_no_auto_pack(request, workdir: Path, parameters: Parameters, cmake_file_api) → NoneView on GitHub : The compiler never automatically packs in CODE_AUTO_PACK.

test_packed() → NoneView on GitHub 

class tests.testing.binaries.sass.test_half.TestFp16FusedMulAddMatcherView on GitHub 

Bases: object

Tests for reprospect.testing.binaries.sass.instruction.half.Fp16FusedMulAddMatcher.

PTX: Final[Path] = PosixPath('/__w/reprospect/reprospect/tests/testing/binaries/sass/assets/hfma2.ptx')

test_any() → NoneView on GitHub 

test_from_ptx(request, workdir: Path, parameters: Parameters) → NoneView on GitHub : Compile the PTX from PTX.

test_individual() → NoneView on GitHub 

test_packed() → NoneView on GitHub 

class tests.testing.binaries.sass.test_half.TestFp16MinMaxMatcherView on GitHub 

Bases: object

Tests for reprospect.testing.binaries.sass.instruction.half.Fp16MinMaxMatcher.

CODE_HMNMX: Final[str] = '#include "cuda_fp16.h"\n\n__global__ void test_h{which}(__half* __restrict__ const out, const __half* __restrict__ const src_a, const __half* __restrict__ const src_b, const unsigned int size)\n{{\n const auto index = blockIdx.x * blockDim.x + threadIdx.x;\n if (index < size)\n out[index] = __h{which}(src_a[index], src_b[index]);\n}}\n'

test_any() → NoneView on GitHub 

test_from_object(request, workdir: Path, parameters: Parameters, which: Literal['min', 'max'], cmake_file_api) → NoneView on GitHub : Inspect the behavior of __hmin and __hmax intrinsics with CODE_HMNMX.

test_max() → NoneView on GitHub 

test_min() → NoneView on GitHub 

class tests.testing.binaries.sass.test_half.TestFp16MulMatcherView on GitHub 

Bases: object

Tests for reprospect.testing.binaries.sass.instruction.half.Fp16MulMatcher.

test_any() → NoneView on GitHub 

test_individual() → NoneView on GitHub 

test_packed() → NoneView on GitHub 

tests.testing.binaries.sass.test_instruction.CODE_ELEMENTWISE_ADD_RESTRICT = '__global__ void elementwise_add_restrict(int* __restrict__ const dst, const int* __restrict__ const src) {\n const auto index = blockIdx.x * blockDim.x + threadIdx.x;\n dst[index] += src[index];\n}\n': Element-wise add with 32-bit int.

tests.testing.binaries.sass.test_instruction.CODE_ELEMENTWISE_ADD_RESTRICT_128_WIDE = '__global__ void elementwise_add_restrict_128_wide(float4* __restrict__ const dst, const float4* __restrict__ const src) {\n const auto index = blockIdx.x * blockDim.x + threadIdx.x;\n const float4& a = src[index];\n const float4& b = dst[index];\n dst[index] = make_float4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w);\n}\n': Element-wise add with 128-bit float4.

tests.testing.binaries.sass.test_instruction.CODE_ELEMENTWISE_ADD_RESTRICT_256_WIDE = 'struct alignas(4 * sizeof(double)) Tester {\n double x, y, z, w;\n};\n\n__global__ void elementwise_add_restrict_256_wide(Tester* __restrict__ const dst, const Tester* __restrict__ const src)\n{\n const auto index = blockIdx.x * blockDim.x + threadIdx.x;\n const Tester& a = src[index];\n const Tester& b = dst[index];\n dst[index] = Tester{.x = a.x + b.x, .y = a.y + b.y, .z = a.z + b.z, .w = a.w + b.w};\n}\n': Element-wise add with 256-bit aligned elements.

class tests.testing.binaries.sass.test_instruction.TestOpcodeModsMatcherView on GitHub 

Bases: object

Tests for reprospect.testing.binaries.sass.instruction.instruction.OpcodeModsMatcher.

test_with_descr()View on GitHub 

test_with_minus_sign()View on GitHub 

test_with_reuse()View on GitHub 

test_with_square_brackets()View on GitHub 

class tests.testing.binaries.sass.test_instruction.TestOpcodeModsWithOperandsMatcherView on GitHub 

Bases: object

Tests for reprospect.testing.binaries.sass.instruction.instruction.OpcodeModsWithOperandsMatcher.

test()View on GitHub 

class tests.testing.binaries.sass.test_instruction.TestReductionMatcherView on GitHub 

Bases: object

Tests for reprospect.testing.binaries.sass.instruction.atomic.ReductionMatcher.

CODE_ADD = '__global__ void add({type}* __restrict__ const dst, const {type}* __restrict__ const src)\n{{\n const auto index = blockIdx.x * blockDim.x + threadIdx.x;\n atomicAdd(&dst[index], src[index]);\n}}\n'

CODE_MAX = '__global__ void max({type}* __restrict__ const dst, const {type}* __restrict__ const src)\n{{\n const auto index = blockIdx.x * blockDim.x + threadIdx.x;\n atomicMax(&dst[index], src[index]);\n}}\n'

CODE_SUB = '__global__ void sub(int* __restrict__ const dst, const int* __restrict__ const src)\n{\n const auto index = blockIdx.x * blockDim.x + threadIdx.x;\n atomicSub(&dst[index], src[index]);\n}\n'

pytestmark = [Mark(name='parametrize', args=('parameters', (Parameters(arch=NVIDIAArch(family=<NVIDIAFamily.TURING: 'TURING'>, compute_capability=ComputeCapability(major=7, minor=5))), Parameters(arch=NVIDIAArch(family=<NVIDIAFamily.AMPERE: 'AMPERE'>, compute_capability=ComputeCapability(major=8, minor=0))), Parameters(arch=NVIDIAArch(family=<NVIDIAFamily.AMPERE: 'AMPERE'>, compute_capability=ComputeCapability(major=8, minor=6))), Parameters(arch=NVIDIAArch(family=<NVIDIAFamily.ADA: 'ADA'>, compute_capability=ComputeCapability(major=8, minor=9))), Parameters(arch=NVIDIAArch(family=<NVIDIAFamily.HOPPER: 'HOPPER'>, compute_capability=ComputeCapability(major=9, minor=0))), Parameters(arch=NVIDIAArch(family=<NVIDIAFamily.BLACKWELL: 'BLACKWELL'>, compute_capability=ComputeCapability(major=10, minor=0))), Parameters(arch=NVIDIAArch(family=<NVIDIAFamily.BLACKWELL: 'BLACKWELL'>, compute_capability=ComputeCapability(major=10, minor=3))), Parameters(arch=NVIDIAArch(family=<NVIDIAFamily.BLACKWELL: 'BLACKWELL'>, compute_capability=ComputeCapability(major=12, minor=0))), Parameters(arch=NVIDIAArch(family=<NVIDIAFamily.BLACKWELL: 'BLACKWELL'>, compute_capability=ComputeCapability(major=12, minor=1))))), kwargs={'ids': <class 'str'>})]

test_add_strong_device_double(request, workdir, parameters: Parameters, cmake_file_api: FileAPI)View on GitHub : Test with CODE_ADD for double.

test_add_strong_device_float(request, workdir, parameters: Parameters, cmake_file_api: FileAPI)View on GitHub : Test with CODE_ADD for float.

test_add_strong_device_int(request, workdir, parameters: Parameters, cmake_file_api: FileAPI)View on GitHub : Test with CODE_ADD for int.

test_add_strong_device_unsigned_int(request, workdir, parameters: Parameters, cmake_file_api: FileAPI)View on GitHub : Test with CODE_ADD for unsigned int.

test_add_strong_device_unsigned_long_long_int(request, workdir, parameters: Parameters, cmake_file_api: FileAPI)View on GitHub : Test with CODE_ADD for unsigned long long int.

test_max_strong_device_int(request, workdir, parameters: Parameters, cmake_file_api: FileAPI)View on GitHub : Test with CODE_MAX for int. The modifier is MAX.S32.

test_max_strong_device_long_long_int(request, workdir, parameters: Parameters, cmake_file_api: FileAPI)View on GitHub : Test with CODE_MAX for long long int. The modifier is MAX.S64.

test_max_strong_device_unsigned_int(request, workdir, parameters: Parameters, cmake_file_api: FileAPI)View on GitHub : Test with CODE_MAX for unsigned int. The modifier is MAX.

test_max_strong_device_unsigned_long_long_int(request, workdir, parameters: Parameters, cmake_file_api: FileAPI)View on GitHub : Test with CODE_MAX for unsigned long long int. The modifier is MAX.64.

test_sub_strong_device(request, workdir, parameters: Parameters, cmake_file_api: FileAPI)View on GitHub : Test with CODE_SUB.

tests.testing.binaries.sass.test_instruction.get_decoder(*, cwd: Path, arch: NVIDIAArch, file: Path, cmake_file_api: FileAPI, **kwargs) → tuple[Decoder, Path]View on GitHub : Compile the code in file for arch and return a reprospect.tools.binaries.sass.decoder.Decoder.

class tests.testing.binaries.sass.test_load.TestLoadConstantMatcherView on GitHub 

Bases: object

Tests for reprospect.testing.binaries.sass.instruction.load.LoadConstantMatcher.

CODE_CONSTANT_ARRAY: Final[str] = '__constant__ {type} data[128];\n__global__ __launch_bounds__(128, 1) void ldc({type}* __restrict__ const out)\n{{\n const auto index = blockIdx.x * blockDim.x + threadIdx.x;\n out[index] = data[index];\n}}\n'

INSTRUCTIONS: Final[dict[str, tuple[LoadConstantMatcher, InstructionMatch]]] = {'LDC R1, c[0x0][0x37c]': (LoadConstantMatcher(pattern=regex.Regex('(?P<opcode>(?:LDC|LDCU)) (?P<operands>(?:R[0-9]+|UR[0-9]+)), (?P<operands>(?:(?:!|\\-\\||\\-|\\~|\\|))?c\\[(?P<bank>0x[0-9]+)\\]\\[(?P<offset>(?:0x[0-9A-Fa-f]+|R(?:Z|\\d+)|UR[0-9]+))\\])', flags=regex.V0)), InstructionMatch(opcode='LDC', modifiers=(), operands=('R1', 'c[0x0][0x37c]'), predicate=None, additional={'bank': ['0x0'], 'offset': ['0x37c']})), 'LDC.64 R6, c[0x1][0x398]': (LoadConstantMatcher(pattern=regex.Regex('(?P<opcode>(?:LDC|LDCU))\\.(?P<modifiers>64) (?P<operands>(?:R[0-9]+|UR[0-9]+)), (?P<operands>(?:(?:!|\\-\\||\\-|\\~|\\|))?c\\[(?P<bank>0x[0-9]+)\\]\\[(?P<offset>(?:0x[0-9A-Fa-f]+|R(?:Z|\\d+)|UR[0-9]+))\\])', flags=regex.V0)), InstructionMatch(opcode='LDC', modifiers=('64',), operands=('R6', 'c[0x1][0x398]'), predicate=None, additional={'bank': ['0x1'], 'offset': ['0x398']})), 'LDCU UR4, c[0x2][0x364]': (LoadConstantMatcher(pattern=regex.Regex('(?P<opcode>LDCU) (?P<operands>UR[0-9]+), (?P<operands>(?:(?:!|\\-\\||\\-|\\~|\\|))?c\\[(?P<bank>0x[0-9]+)\\]\\[(?P<offset>(?:0x[0-9A-Fa-f]+|R(?:Z|\\d+)|UR[0-9]+))\\])', flags=regex.V0)), InstructionMatch(opcode='LDCU', modifiers=(), operands=('UR4', 'c[0x2][0x364]'), predicate=None, additional={'bank': ['0x2'], 'offset': ['0x364']})), 'LDCU UR4, c[0x3][UR0]': (LoadConstantMatcher(pattern=regex.Regex('(?P<opcode>LDCU) (?P<operands>UR[0-9]+), (?P<operands>(?:(?:!|\\-\\||\\-|\\~|\\|))?c\\[(?P<bank>0x[0-9]+)\\]\\[(?P<offset>(?:0x[0-9A-Fa-f]+|R(?:Z|\\d+)|UR[0-9]+))\\])', flags=regex.V0)), InstructionMatch(opcode='LDCU', modifiers=(), operands=('UR4', 'c[0x3][UR0]'), predicate=None, additional={'bank': ['0x3'], 'offset': ['UR0']})), 'LDCU.64 UR6, c[0x3][0x358]': (LoadConstantMatcher(pattern=regex.Regex('(?P<opcode>LDCU)\\.(?P<modifiers>64) (?P<operands>UR[0-9]+), (?P<operands>(?:(?:!|\\-\\||\\-|\\~|\\|))?c\\[(?P<bank>0x[0-9]+)\\]\\[(?P<offset>(?:0x[0-9A-Fa-f]+|R(?:Z|\\d+)|UR[0-9]+))\\])', flags=regex.V0)), InstructionMatch(opcode='LDCU', modifiers=('64',), operands=('UR6', 'c[0x3][0x358]'), predicate=None, additional={'bank': ['0x3'], 'offset': ['0x358']}))}: Zoo of real SASS instructions.

test(instruction: str, matcher: LoadConstantMatcher, expected: InstructionMatch) → NoneView on GitHub 

test_array_of_64bit_elements(request, workdir: Path, parameters: Parameters, cmake_file_api: FileAPI) → NoneView on GitHub : Loads of size 64 with CODE_CONSTANT_ARRAY.

class tests.testing.binaries.sass.test_load.TestLoadMatcherView on GitHub 

Bases: object

Tests for reprospect.testing.binaries.sass.instruction.load.LoadMatcher and reprospect.testing.binaries.sass.instruction.load.LoadGlobalMatcher.

CODE_ELEMENTWISE_ADD = '__global__ void elementwise_add(int* const dst, const int* const src) {\n const auto index = blockIdx.x * blockDim.x + threadIdx.x;\n dst[index] += src[index];\n}\n'

CODE_ELEMENTWISE_ADD_LDG = '__global__ void elementwise_add_ldg(int* const dst, const int* const src) {\n const auto index = blockIdx.x * blockDim.x + threadIdx.x;\n dst[index] += __ldg(&src[index]);\n}\n'

CODE_EXTEND = '#include <cstdint>\n\n__global__ void extend({dst}* {restrict} const dst, {src}* {restrict} const src, const unsigned int size)\n{{\n const auto index = blockIdx.x * blockDim.x + threadIdx.x;\n if (index < size) dst[index] = src[index];\n}}\n'

test() → NoneView on GitHub 

test_constant(request, workdir, parameters: Parameters, cmake_file_api: FileAPI) → NoneView on GitHub : If src is declared const __restrict__, the compiler is able to use the .CONSTANT modifier. Otherwise, we need to explicitly use __ldg to end up using .CONSTANT.

test_elementwise_add_restrict(request, workdir, parameters: Parameters, cmake_file_api: FileAPI) → NoneView on GitHub : Test loads with tests.testing.binaries.sass.test_instruction.CODE_ELEMENTWISE_ADD_RESTRICT.

test_elementwise_add_restrict_128_wide(request, workdir, parameters: Parameters, cmake_file_api: FileAPI) → NoneView on GitHub : Test 128-bit wide loads with tests.testing.binaries.sass.test_instruction.CODE_ELEMENTWISE_ADD_RESTRICT_128_WIDE.

test_elementwise_add_restrict_256_wide(request, workdir, parameters: Parameters, cmake_file_api: FileAPI) → NoneView on GitHub : Test 256-bit wide loads with tests.testing.binaries.sass.test_instruction.CODE_ELEMENTWISE_ADD_RESTRICT_256_WIDE.

test_sign_extend_s16(request, workdir: Path, parameters: Parameters, cmake_file_api: FileAPI, cmake_cuda_compiler: CMakeToolchainCompiler) → NoneView on GitHub 

Check when CODE_EXTEND leads to sign extension.

Uses reprospect.testing.features.Memory.sign_extension().

test_zero_extend_u16(request, workdir: Path, parameters: Parameters, cmake_file_api: FileAPI) → NoneView on GitHub : Use CODE_EXTEND to enforce zero extension.

class tests.testing.binaries.sass.test_store.TestStoreMatcherView on GitHub 

Bases: object

Tests for reprospect.testing.binaries.sass.instruction.store.StoreMatcher and reprospect.testing.binaries.sass.instruction.store.StoreGlobalMatcher.

test() → NoneView on GitHub 

test_elementwise_add_restrict(request, workdir, parameters: Parameters, cmake_file_api: FileAPI) → NoneView on GitHub : Test store with tests.testing.binaries.sass.test_instruction.CODE_ELEMENTWISE_ADD_RESTRICT.

test_elementwise_add_restrict_128_wide(request, workdir, parameters: Parameters, cmake_file_api: FileAPI) → NoneView on GitHub : Test 128-bit wide store with tests.testing.binaries.sass.test_instruction.CODE_ELEMENTWISE_ADD_RESTRICT_128_WIDE.

test_elementwise_add_restrict_256_wide(request, workdir, parameters: Parameters, cmake_file_api: FileAPI) → NoneView on GitHub : Test 256-bit wide store with tests.testing.binaries.sass.test_instruction.CODE_ELEMENTWISE_ADD_RESTRICT_256_WIDE.