Instruction

class tests.testing.binaries.sass.instruction.test_any.TestAnyMatcherView on GitHub

Bases: object

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

INSTRUCTIONS: Final[dict[str, InstructionMatch]] = {'@!P3 LDG.E.64.CONSTANT R6, desc[UR16][R26.64]': InstructionMatch(opcode='LDG', modifiers=('E', '64', 'CONSTANT'), operands=('R6', 'desc[UR16][R26.64]'), predicate='@!P3', additional=None), '@!PT LDS RZ, [RZ]': InstructionMatch(opcode='LDS', modifiers=(), operands=('RZ', '[RZ]'), predicate='@!PT', additional=None), '@!UP0 UIMAD UR7, UR10, 0xc, URZ': InstructionMatch(opcode='UIMAD', modifiers=(), operands=('UR7', 'UR10', '0xc', 'URZ'), predicate='@!UP0', additional=None), '@UP0 LDCU.64 UR12, c[0x0][0x388]': InstructionMatch(opcode='LDCU', modifiers=('64',), operands=('UR12', 'c[0x0][0x388]'), predicate='@UP0', additional=None), 'ATOMG.E.ADD.STRONG.GPU PT, R4, desc[UR16][R4.64], R7': InstructionMatch(opcode='ATOMG', modifiers=('E', 'ADD', 'STRONG', 'GPU'), operands=('PT', 'R4', 'desc[UR16][R4.64]', 'R7'), predicate=None, additional=None), 'BAR.SYNC.DEFER_BLOCKING 0x0': InstructionMatch(opcode='BAR', modifiers=('SYNC', 'DEFER_BLOCKING'), operands=('0x0',), predicate=None, additional=None), 'BRA 0x240': InstructionMatch(opcode='BRA', modifiers=(), operands=('0x240',), predicate=None, additional=None), 'DADD R12, |R12|, |R14|': InstructionMatch(opcode='DADD', modifiers=(), operands=('R12', '|R12|', '|R14|'), predicate=None, additional=None), 'EXIT': InstructionMatch(opcode='EXIT', modifiers=(), operands=(), predicate=None, additional=None), 'F2FP.BF16.PACK_AB R27, R6, R27': InstructionMatch(opcode='F2FP', modifiers=('BF16', 'PACK_AB'), operands=('R27', 'R6', 'R27'), predicate=None, additional=None), 'FADD.FTZ.RN R0, R1, R2': InstructionMatch(opcode='FADD', modifiers=('FTZ', 'RN'), operands=('R0', 'R1', 'R2'), predicate=None, additional=None), 'FMUL R6, R27.reuse, 0.044714998453855514526': InstructionMatch(opcode='FMUL', modifiers=(), operands=('R6', 'R27.reuse', '0.044714998453855514526'), predicate=None, additional=None), 'FMUL.FTZ R2, R2, R3': InstructionMatch(opcode='FMUL', modifiers=('FTZ',), operands=('R2', 'R2', 'R3'), predicate=None, additional=None), 'HFMA2 R7, -RZ, RZ, 0, 5.9604644775390625e-08': InstructionMatch(opcode='HFMA2', modifiers=(), operands=('R7', '-RZ', 'RZ', '0', '5.9604644775390625e-08'), predicate=None, additional=None), 'HMUL2 R0, R2.H0_H0, R3.H0_H0': InstructionMatch(opcode='HMUL2', modifiers=(), operands=('R0', 'R2.H0_H0', 'R3.H0_H0'), predicate=None, additional=None), 'IADD3 R2, R2, 0x4, RZ': InstructionMatch(opcode='IADD3', modifiers=(), operands=('R2', 'R2', '0x4', 'RZ'), predicate=None, additional=None), 'IMAD.MOV.U32 R4, R4, R5, R6': InstructionMatch(opcode='IMAD', modifiers=('MOV', 'U32'), operands=('R4', 'R4', 'R5', 'R6'), predicate=None, additional=None), 'ISETP.NE.AND P0, PT, R1, RZ, PT': InstructionMatch(opcode='ISETP', modifiers=('NE', 'AND'), operands=('P0', 'PT', 'R1', 'RZ', 'PT'), predicate=None, additional=None), 'LDG.E R0, desc[UR12][R18.64]': InstructionMatch(opcode='LDG', modifiers=('E',), operands=('R0', 'desc[UR12][R18.64]'), predicate=None, additional=None), 'LDG.E.SYS R4, [R2]': InstructionMatch(opcode='LDG', modifiers=('E', 'SYS'), operands=('R4', '[R2]'), predicate=None, additional=None), 'LDGSTS.E.LTC128B [R211+-0x4000], desc[UR24][R182.64+0x20]': InstructionMatch(opcode='LDGSTS', modifiers=('E', 'LTC128B'), operands=('[R211+-0x4000]', 'desc[UR24][R182.64+0x20]'), predicate=None, additional=None), 'LDS.64 R12, [UR7+0x8]': InstructionMatch(opcode='LDS', modifiers=('64',), operands=('R12', '[UR7+0x8]'), predicate=None, additional=None), 'LEA.HI.X R145, R140, R134.reuse, R135, 0x2, P3': InstructionMatch(opcode='LEA', modifiers=('HI', 'X'), operands=('R145', 'R140', 'R134.reuse', 'R135', '0x2', 'P3'), predicate=None, additional=None), 'LOP3.LUT P0, RZ, R8, R94, RZ, 0xfc, !PT': InstructionMatch(opcode='LOP3', modifiers=('LUT',), operands=('P0', 'RZ', 'R8', 'R94', 'RZ', '0xfc', '!PT'), predicate=None, additional=None), 'MEMBAR.SC.GPU': InstructionMatch(opcode='MEMBAR', modifiers=('SC', 'GPU'), operands=(), predicate=None, additional=None), 'MOV R8, c[0x0][0x140]': InstructionMatch(opcode='MOV', modifiers=(), operands=('R8', 'c[0x0][0x140]'), predicate=None, additional=None), 'NOP': InstructionMatch(opcode='NOP', modifiers=(), operands=(), predicate=None, additional=None), 'RET.REL.NODEC R4 0x0': InstructionMatch(opcode='RET', modifiers=('REL', 'NODEC'), operands=('R4', '0x0'), predicate=None, additional=None), 'S2R R0, SR_CTAID.X': InstructionMatch(opcode='S2R', modifiers=(), operands=('R0', 'SR_CTAID.X'), predicate=None, additional=None), 'STG.E [R2], R4': InstructionMatch(opcode='STG', modifiers=('E',), operands=('[R2]', 'R4'), predicate=None, additional=None), 'STG.E.64.SYS [R2.64+UR4], R4': InstructionMatch(opcode='STG', modifiers=('E', '64', 'SYS'), operands=('[R2.64+UR4]', 'R4'), predicate=None, additional=None), 'STL [R1], R2': InstructionMatch(opcode='STL', modifiers=(), operands=('[R1]', 'R2'), predicate=None, additional=None), 'STS.64 [R9.X8], R2': InstructionMatch(opcode='STS', modifiers=('64',), operands=('[R9.X8]', 'R2'), predicate=None, additional=None), 'UIADD3 UR5, UPT, UPT, -UR4, UR9, URZ': InstructionMatch(opcode='UIADD3', modifiers=(), operands=('UR5', 'UPT', 'UPT', '-UR4', 'UR9', 'URZ'), predicate=None, additional=None)}

Zoo of real SASS instructions.

MATCHER: Final[AnyMatcher] = AnyMatcher(pattern=regex.Regex('(?:(?P<predicate>@!?U?P(?:T|[0-9]+)))?\\s*(?P<opcode>[A-Z0-9]+)(?:\\.(?P<modifiers>[A-Z0-9_]+))*\\s*(?:(?P<operands>[\\w!\\.\\[\\]\\+\\-\\|~]+)(?:,?\\s*(?P<operands>[\\w!\\.\\[\\]\\+\\-\\|~]+))*)?', flags=regex.V0))

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

test_no_match() → NoneView on GitHub

test_pattern() → NoneView on GitHub

class tests.testing.binaries.sass.instruction.test_atomic.TestAtomicMatcherView on GitHub

Bases: object

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

CODE_ADD_BASED_ON_CAS = '__global__ void cas({type}* __restrict__ const dst, const {type}* __restrict__ const src)\n{{\n    static_assert(sizeof({type}) == {size});\n\n    const auto index = blockIdx.x * blockDim.x + threadIdx.x;\n\n    {integer}* const dest = reinterpret_cast<{integer}*>(dst + index);\n\n    {integer} old = *dest;\n    {integer} assumed;\n\n    do {{\n        assumed = old;\n        {type} new_val = assumed + src[index];\n        old = atomicCAS(\n            dest,\n            assumed,\n            reinterpret_cast<{integer}&>(new_val)\n        );\n    }} while (old != assumed);\n}}\n'

CODE_ADD_BASED_ON_CAS_128 = 'struct alignas(2 * sizeof(double)) My128Struct\n{\n    double x, y;\n\n    __host__ __device__ friend My128Struct operator+(const My128Struct& a, const My128Struct& b)\n    {\n        return My128Struct{\n            .x = a.x + b.x,\n            .y = a.y + b.y\n        };\n    }\n\n    auto operator<=>(const My128Struct&) const = default;\n};\n\n__global__ void cas(My128Struct* __restrict__ const dst, const My128Struct* __restrict__ const src)\n{\n    static_assert(sizeof (My128Struct) == 16);\n    static_assert(alignof(My128Struct) == 16);\n\n    static_assert(std::is_trivially_copyable_v<My128Struct>);\n\n    const auto index = blockIdx.x * blockDim.x + threadIdx.x;\n\n    auto* const dest = dst + index;\n\n    My128Struct old = *dest;\n    My128Struct assumed;\n\n    do {\n        assumed = old;\n        My128Struct new_val = assumed + src[index];\n        old = atomicCAS(\n            dest,\n            assumed,\n            new_val\n        );\n    } while (old != assumed);\n}\n'

CODE_ADD_RELAXED_BLOCK = '#include "cuda/atomic"\n\n__global__ void add({type}* __restrict__ const dst, const {type}* __restrict__ const src)\n{{\n    const auto index = blockIdx.x * blockDim.x + threadIdx.x;\n\n    cuda::atomic_ref<{type}, cuda::thread_scope_block> ref(dst[index]);\n    ref.fetch_add(src[index], cuda::memory_order_relaxed);\n}}\n'

CODE_COMPARE_EXCHANGE = '#include "cuda/atomic"\n\n__global__ void compare_exchange({type}* __restrict__ dst, const {type}* __restrict__ src)\n{{\n    const auto index = blockIdx.x * blockDim.x + threadIdx.x;\n    cuda::atomic_ref<{type}, cuda::thread_scope_{scope}> ref(dst[index]);\n\n    {type} old_val = ref.load(cuda::memory_order_relaxed);\n    {type} new_val;\n\n    do {{\n        new_val = min(old_val, src[index]);\n    }} while (!ref.compare_exchange_{consistency}(old_val, new_val,\n                                       cuda::memory_order_acquire,\n                                       cuda::memory_order_acquire));\n}}\n'

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

CODE_EXCH_DEVICE_PTR: Final[str] = '__device__ __constant__ int32_t* ptr;\n\n__global__ void atomic_exch_kernel() {\n    atomicExch(ptr, 0);\n}\n'

CODE_MIN = '#include "cuda/atomic"\n\n__global__ void add({type}* __restrict__ const dst, const {type}* __restrict__ const src)\n{{\n    const auto index = blockIdx.x * blockDim.x + threadIdx.x;\n    cuda::atomic_ref<{type}, cuda::thread_scope_{scope}> ref(dst[index]);\n    ref.fetch_min(src[index], cuda::memory_order_relaxed);\n}}\n'

static assert_atomicCAS_ptx(*, output: Path, matcher: AtomicMatcher, cuda_compiler: CMakeToolchainCompiler) → NoneView on GitHub

Check for the expected PTX code in output given the reprospect.testing.binaries.sass.instruction.atomic.ThreadScope.

According to:

• https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html#legacy-atomic-functions

it should generate reprospect.testing.binaries.sass.instruction.atomic.ThreadScope.DEVICE instructions. However, as noted in:

• https://github.com/uliegecsm/reprospect/issues/556

• https://forums.developer.nvidia.com/t/atomiccas-cuda-13-2-unexpectedly-generates-system-scope-atomic-instruction-starting-from-cc100/365449

it may unexpectedly generate reprospect.testing.binaries.sass.instruction.atomic.ThreadScope.SYSTEM instructions.

Note

According to https://docs.nvidia.com/cuda/parallel-thread-execution/#state-spaces-state-spaces-tab, .global is global memory shared by all threads.

Note

According to https://docs.nvidia.com/cuda/parallel-thread-execution/#id682, .sys is the set of all threads in the current program, including all kernel grids invoked by the host program on all compute devices, and all threads constituting the host program itself.

static get_atomicCAS_thread_scope(*, size: int, arch: NVIDIAArch, cuda_compiler: CMakeToolchainCompiler) → ThreadScopeView on GitHub

Get the expected reprospect.testing.binaries.sass.instruction.atomic.ThreadScope for atomicCAS.

static match_one(*, decoder, **kwargs) → tuple[AtomicMatcher, Instruction, InstructionMatch]View on GitHub

Match exactly one instruction.

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_relaxed_block_double(request, workdir, parameters: Parameters, cmake_file_api: FileAPI)View on GitHub

Similar to test_add_relaxed_block_int() for double.

test_add_relaxed_block_float(request, workdir, parameters: Parameters, cmake_file_api: FileAPI)View on GitHub

Similar to test_add_relaxed_block_int() for float.

test_add_relaxed_block_int(request, workdir, parameters: Parameters, cmake_file_api: FileAPI)View on GitHub

As of CUDA 13.0.0, the generated code still applies the .STRONG modifier, regardless of the .relaxed qualifier shown in the PTX.

test_add_relaxed_block_unsigned_long_long_int(request, workdir, parameters: Parameters, cmake_file_api: FileAPI)View on GitHub

Similar to test_add_relaxed_block_int() for unsigned long long int.

test_atomicCAS(request, workdir, word, parameters: Parameters, cmake_file_api: FileAPI, cmake_cuda_compiler: CMakeToolchainCompiler)View on GitHub

Test with CODE_ADD_BASED_ON_CAS.

test_atomicCAS_128(request, workdir, parameters: Parameters, cmake_file_api: FileAPI, cmake_cuda_compiler: CMakeToolchainCompiler)View on GitHub

Supported from compute capability 9.x.

test_compare_exchange_system(request, consistency: str, workdir, parameters: Parameters, cmake_file_api: FileAPI)View on GitHub

Test with CODE_COMPARE_EXCHANGE for unsigned long long int, reprospect.testing.binaries.sass.instruction.atomic.ThreadScope.SYSTEM scope.

Note

Both weak and strong lead to a STRONG consistency.

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

This test demonstrates that while nvcc emits an ATOMG instruction for an atomic exchange using a device pointer marked with __constant__, clang (as of 21.1.5) is not able to infer that the referenced memory resides in global memory and therefore falls back emitting a generic ATOM instruction.

nvcc appears to generate better code in this case: because the device pointer is declared as __constant__, the compiler can reasonably assume that it cannot point to local or shared memory, and thus must refer to global memory. This allows nvcc to use the more specific global-memory atomic instruction.

test_exch_strong_device_float(request, workdir, parameters: Parameters, cmake_file_api: FileAPI)View on GitHub

Test with CODE_EXCH for float.

test_exch_strong_device_int(request, workdir, parameters: Parameters, cmake_file_api: FileAPI)View on GitHub

Test with CODE_EXCH for int.

test_exch_strong_device_unsigned_long_long_int(request, workdir, parameters: Parameters, cmake_file_api: FileAPI)View on GitHub

Test with CODE_EXCH for unsigned long long int.

test_min_relaxed_device_int(request, workdir, parameters: Parameters, cmake_file_api: FileAPI)View on GitHub

Test with CODE_MIN for int and reprospect.testing.binaries.sass.instruction.atomic.ThreadScope.DEVICE scope.

test_min_relaxed_device_long_long_int(request, workdir, parameters: Parameters, cmake_file_api: FileAPI)View on GitHub

Test with CODE_MIN for long long int and reprospect.testing.binaries.sass.instruction.atomic.ThreadScope.DEVICE scope.

test_min_relaxed_device_unsigned_long_long_int(request, workdir, parameters: Parameters, cmake_file_api: FileAPI)View on GitHub

Test with CODE_MIN for unsigned long long int and reprospect.testing.binaries.sass.instruction.atomic.ThreadScope.DEVICE scope.

test_min_relaxed_system_int(request, workdir, parameters: Parameters, cmake_file_api: FileAPI)View on GitHub

Test with CODE_MIN for int and reprospect.testing.binaries.sass.instruction.atomic.ThreadScope.SYSTEM scope.

class tests.testing.binaries.sass.instruction.test_branch.TestBranchMatcherView on GitHub

Bases: object

Tests for reprospect.testing.binaries.sass.instruction.branch.BranchMatcher.

INSTRUCTIONS: Final[dict[str, InstructionMatch]] = {'@!UP0 BRA 0xc': InstructionMatch(opcode='BRA', modifiers=(), operands=('0xc',), predicate='@!UP0', additional=None), '@PT BRA 0xf0': InstructionMatch(opcode='BRA', modifiers=(), operands=('0xf0',), predicate='@PT', additional=None), '@UPT BRA 0x170': InstructionMatch(opcode='BRA', modifiers=(), operands=('0x170',), predicate='@UPT', additional=None), 'BRA 0x240': InstructionMatch(opcode='BRA', modifiers=(), operands=('0x240',), predicate=None, additional=None)}

Zoo of real BRA instructions.

MATCHER: Final[BranchMatcher] = BranchMatcher(pattern=regex.Regex('(?:(?P<predicate>@!?U?P(?:T|[0-9]+)))?\\s*(?P<opcode>BRA)\\s*(?P<operands>0x[0-9A-Fa-f]+)', flags=regex.V0))

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

test_no_match() → NoneView on GitHub

test_pattern() → NoneView on GitHub

test_with_predicate() → NoneView on GitHub

Verify that the matcher only matches instructions containing the exact predicate. Instructions without the predicate or with a different predicate should not match.

class tests.testing.binaries.sass.instruction.test_address.TestAddressMatcherView on GitHub

Bases: object

Tests for reprospect.testing.binaries.sass.instruction.address.AddressMatcher.

test(parameters: Parameters) → NoneView on GitHub

Ensure that reprospect.testing.binaries.sass.instruction.address.AddressMatcher.build_pattern() supports all tests.parameters.PARAMETERS.

test_build_pattern() → NoneView on GitHub

test_desc_reg64_address() → NoneView on GitHub

test_extend_bits_address_turing75(request, workdir: Path, cmake_file_api: FileAPI) → NoneView on GitHub

Generate a SASS instruction such as:

ST.E.STRONG.SYS [R4.U32+UR4+0x4], R3

and match it.

test_global_reg64_address_turing75() → NoneView on GitHub

test_reg64_address() → NoneView on GitHub

test_reg_address() → NoneView on GitHub

test_reg_stride_address() → NoneView on GitHub

test_shared_offset_address_ampere86() → NoneView on GitHub

test_shared_offset_address_volta70() → NoneView on GitHub

test_shared_rz_address() → NoneView on GitHub

test_shared_stride_16() → NoneView on GitHub

test_shared_ur() → NoneView on GitHub

class tests.testing.binaries.sass.instruction.test_constant.TestConstantMatcherView on GitHub

Bases: object

Tests for reprospect.testing.binaries.sass.instruction.constant.ConstantMatcher.

CONSTANTS: Final[dict[str, ConstantMatch]] = {'-c[0x0][0x18c]': ConstantMatch(bank='0x0', offset='0x18c', math=<MathModifier.NEG: '-'>), 'c[0x0][0x37c]': ConstantMatch(bank='0x0', offset='0x37c', math=None), 'c[0x0][R9]': ConstantMatch(bank='0x0', offset='R9', math=None), 'c[0x0][RZ]': ConstantMatch(bank='0x0', offset='RZ', math=None), 'c[0x0][UR456]': ConstantMatch(bank='0x0', offset='UR456', math=None)}

MATCHER: Final[ConstantMatcher] = ConstantMatcher(bank=None, offset=None, math=None, pattern=regex.Regex('(?:(?P<modifier_math>(?:!|\\-\\||\\-|\\~|\\|)))?c\\[(?P<bank>0x[0-9]+)\\]\\[(?P<offset>(?:0x[0-9A-Fa-f]+|R(?:Z|\\d+)|UR[0-9]+))\\]', flags=regex.V0))

test_any(constant: str, expected: ConstantMatch) → NoneView on GitHub

test_build_pattern() → NoneView on GitHub

test_constant() → NoneView on GitHub

test_constant_neg() → NoneView on GitHub

test_constant_with_reg() → NoneView on GitHub

test_constant_with_rz() → NoneView on GitHub

test_constant_with_ureg() → NoneView on GitHub

class tests.testing.binaries.sass.instruction.test_cublas.TestCuBLASView on GitHub

Bases: object

test_operands(workdir: Path, parameters: Parameters, cmake_file_api: FileAPI) → NoneView on GitHub

For each supported architecture, pick one cubin from cuBLAS. Then, loop over all instructions of each function, and match each operand.

tests.testing.binaries.sass.instruction.test_cublas.fixit(instruction: Instruction) → InstructionView on GitHub

Fix the few instructions that do not conform to the expected format.

Before	After
]␣[	][
␣,	,

Here are a few examples:

• HADD2.F32 R10, -RZ, c[0x0]␣[0x160].H0_H0 becomes HADD2.F32 R10, -RZ, c[0x0][0x160].H0_H0.

• FSETP.GEU.AND P1, PT, |R151|, +INF␣, PT becomes FSETP.GEU.AND P1, PT, |R151|, +INF, PT.

These cases are rare (a handful out of millions of instructions), so the search acts as a fast-path guard before the substitution.

class tests.testing.binaries.sass.instruction.test_fp32add.TestFp32AddMatcherView on GitHub

Bases: object

Tests for reprospect.testing.binaries.sass.instruction.floating.Fp32AddMatcher.

test() → NoneView on GitHub

test_elementwise_add_restrict_wide(request, workdir, parameters: Parameters, cmake_file_api: FileAPI)View on GitHub

Test with tests.testing.binaries.sass.test_instruction.CODE_ELEMENTWISE_ADD_RESTRICT_128_WIDE.

There will be 4 FADD instructions because of the float4.

class tests.testing.binaries.sass.instruction.test_fp64add.TestFp64AddMatcherView on GitHub

Bases: object

Tests for reprospect.testing.binaries.sass.instruction.floating.Fp64AddMatcher.

CODE_FP64_ADD: Final[str] = '__global__ void fp64_add(double* __restrict__ const dst, const double* __restrict__ const src)\n{\n    const auto index = blockIdx.x * blockDim.x + threadIdx.x;\n    dst[index] += src[index];\n}\n'

test() → NoneView on GitHub

test_from_compiled(request, workdir, parameters: Parameters, cmake_file_api: FileAPI)View on GitHub

Test with CODE_FP64_ADD.

class tests.testing.binaries.sass.instruction.test_immediate.TestImmediateView on GitHub

Bases: object

Tests for reprospect.testing.binaries.sass.instruction.immediate.Immediate.

IMMEDIATES: Final[tuple[str, ...]] = ('3.1400001049041748047', '1.00000000000000000000e+10', '0.5', '100', '-0.5', '+QNAN', '0x7fffffff', '+INF', '-QNAN', '-INF')

Zoo of real floating-point immediates.

test_any(immediate: str) → NoneView on GitHub

test_floating() → NoneView on GitHub

test_floating_or_limit() → NoneView on GitHub

test_inf() → NoneView on GitHub

test_qnan() → NoneView on GitHub

class tests.testing.binaries.sass.instruction.test_integer.TestIntAdd3MatcherView on GitHub

Bases: object

Tests for reprospect.testing.binaries.sass.instruction.integer.IntAdd3Matcher.

test_as_of_blackwell() → NoneView on GitHub

test_before_blackwell() → NoneView on GitHub

class tests.testing.binaries.sass.instruction.test_integer.TestIntAddMatcherView on GitHub

Bases: object

Tests for reprospect.testing.binaries.sass.instruction.integer.IntAddMatcher.

test() → NoneView on GitHub

class tests.testing.binaries.sass.instruction.test_integer.TestLEAMatcherView on GitHub

Bases: object

Tests for reprospect.testing.binaries.sass.instruction.integer.LEAMatcher.

test() → NoneView on GitHub

test_shift() → NoneView on GitHub

Enforce the shift.

class tests.testing.binaries.sass.instruction.test_operand.TestOperandView on GitHub

Bases: object

Tests for reprospect.testing.binaries.sass.instruction.operand.Operand.

OPERANDS: Final[tuple[str, ...]] = ('SR_TID.X', 'R126', 'UR24', 'R125.reuse', '-R128', '~URZ', '|R165|', '!PT', '[R135]', '[R155+0x1200]', '[R192.X8+0x2000]', 'desc[UR22][R10.64]', 'desc[UR22][R16.64+0x80]', 'c[0x0][0x358]', '100', '0xffffffe0', '-0xfe', '+INF')

Zoo of real operands.

test_any(opnd: str) → NoneView on GitHub

test_mod() → NoneView on GitHub

class tests.testing.binaries.sass.instruction.test_register.TestRegisterView on GitHub

Bases: object

Tests for reprospect.testing.binaries.sass.instruction.register.Register.

test_reg_vs_ureg() → NoneView on GitHub

Ensure that reprospect.testing.binaries.sass.instruction.register.Register.REG does not match what reprospect.testing.binaries.sass.instruction.register.Register.UREG matches (and vice versa).

class tests.testing.binaries.sass.instruction.test_register.TestRegisterMatcherView on GitHub

Bases: object

Tests for reprospect.testing.binaries.sass.instruction.register.RegisterMatcher.

MATCHER: Final[RegisterMatcher] = RegisterMatcher(rtype=None, special=None, index=None, reuse=None, math=None, pattern=regex.Regex('(?:(?P<modifier_math>(?:!|\\-\\||\\-|\\~|\\|)))?(?P<rtype>(?:R|UR|P|UP))(?:(?P<special>(?:Z|T))|(?P<index>\\d+))(?:\\.(?P<reuse>reuse))?', flags=regex.V0))

REGISTERS: Final[dict[str, RegisterMatch]] = {'!P0': RegisterMatch(rtype=<RegisterType.PRED: 'P'>, index=0, reuse=False, math=<MathModifier.NOT: '!'>), '!R42': RegisterMatch(rtype=<RegisterType.GPR: 'R'>, index=42, reuse=False, math=<MathModifier.NOT: '!'>), '-R42': RegisterMatch(rtype=<RegisterType.GPR: 'R'>, index=42, reuse=False, math=<MathModifier.NEG: '-'>), '-|R42|': RegisterMatch(rtype=<RegisterType.GPR: 'R'>, index=42, reuse=False, math=<MathModifier.NEG_ABS: '-|'>), 'P3': RegisterMatch(rtype=<RegisterType.PRED: 'P'>, index=3, reuse=False, math=None), 'PT': RegisterMatch(rtype=<RegisterType.PRED: 'P'>, index=-1, reuse=False, math=None), 'R42': RegisterMatch(rtype=<RegisterType.GPR: 'R'>, index=42, reuse=False, math=None), 'R42.reuse': RegisterMatch(rtype=<RegisterType.GPR: 'R'>, index=42, reuse=True, math=None), 'RZ': RegisterMatch(rtype=<RegisterType.GPR: 'R'>, index=-1, reuse=False, math=None), 'UP3': RegisterMatch(rtype=<RegisterType.UPRED: 'UP'>, index=3, reuse=False, math=None), 'UPT': RegisterMatch(rtype=<RegisterType.UPRED: 'UP'>, index=-1, reuse=False, math=None), 'UR42': RegisterMatch(rtype=<RegisterType.UGPR: 'UR'>, index=42, reuse=False, math=None), '|R42|': RegisterMatch(rtype=<RegisterType.GPR: 'R'>, index=42, reuse=False, math=<MathModifier.ABS: '|'>), '~R42': RegisterMatch(rtype=<RegisterType.GPR: 'R'>, index=42, reuse=False, math=<MathModifier.INV: '~'>)}

test_any(register, expected) → NoneView on GitHub

test_not_a_reg() → NoneView on GitHub

test_pred() → NoneView on GitHub

test_pred_not() → NoneView on GitHub

test_predt() → NoneView on GitHub

test_reg() → NoneView on GitHub

test_reg_no_math() → NoneView on GitHub

test_reg_reuse() → NoneView on GitHub

test_regz() → NoneView on GitHub

test_upred() → NoneView on GitHub

test_upredt() → NoneView on GitHub

test_ureg() → NoneView on GitHub