Instruction

class tests.test.sass.instruction.test_address.TestAddressMatcherView on GitHub

Bases: object

Tests for reprospect.test.sass.instruction.address.AddressMatcher.

test(parameters: Parameters) → NoneView on GitHub

Ensure that reprospect.test.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_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() → 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.test.sass.instruction.test_constant.TestConstantMatcherView on GitHub

Bases: object

Tests for reprospect.test.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.test.sass.instruction.test_fp32add.TestFp32AddMatcherView on GitHub

Bases: object

Tests for reprospect.test.sass.instruction.Fp32AddMatcher.

test() → NoneView on GitHub

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

Test with tests.test.sass.test_instruction.CODE_ELEMENTWISE_ADD_RESTRICT_128_WIDE.

There will be 4 FADD instructions because of the float4.

class tests.test.sass.instruction.test_fp64add.TestFp64AddMatcherView on GitHub

Bases: object

Tests for reprospect.test.sass.instruction.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.test.sass.instruction.test_half.TestFp16AddMatcherView on GitHub

Bases: object

Tests for reprospect.test.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.test.sass.instruction.test_half.TestFp16FusedMulAddMatcherView on GitHub

Bases: object

Tests for reprospect.test.sass.instruction.half.Fp16FusedMulAddMatcher.

PTX: Final[Path] = PosixPath('/__w/reprospect/reprospect/tests/test/sass/instruction/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.test.sass.instruction.test_half.TestFp16MinMaxMatcherView on GitHub

Bases: object

Tests for reprospect.test.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.test.sass.instruction.test_half.TestFp16MulMatcherView on GitHub

Bases: object

Tests for reprospect.test.sass.instruction.half.Fp16MulMatcher.

test_any() → NoneView on GitHub

test_individual() → NoneView on GitHub

test_packed() → NoneView on GitHub

class tests.test.sass.instruction.test_immediate.TestImmediateView on GitHub

Bases: object

Tests for reprospect.test.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.test.sass.instruction.test_integer.TestIntAdd3MatcherView on GitHub

Bases: object

Tests for reprospect.test.sass.instruction.integer.IntAdd3Matcher.

test_as_of_blackwell() → NoneView on GitHub

test_before_blackwell() → NoneView on GitHub

class tests.test.sass.instruction.test_integer.TestIntAddMatcherView on GitHub

Bases: object

Tests for reprospect.test.sass.instruction.integer.IntAddMatcher.

test() → NoneView on GitHub

class tests.test.sass.instruction.test_integer.TestLEAMatcherView on GitHub

Bases: object

Tests for reprospect.test.sass.instruction.integer.LEAMatcher.

test() → NoneView on GitHub

test_shift() → NoneView on GitHub

Enforce the shift.

class tests.test.sass.instruction.test_operand.TestOperandView on GitHub

Bases: object

Tests for reprospect.test.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.test.sass.instruction.test_register.TestRegisterView on GitHub

Bases: object

Tests for reprospect.test.sass.instruction.register.Register.

test_reg_vs_ureg() → NoneView on GitHub

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

class tests.test.sass.instruction.test_register.TestRegisterMatcherView on GitHub

Bases: object

Tests for reprospect.test.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