SPIRV-Runner crash after signature/constant representation change (PR-3043)

[kballeda]

Describe the bug

SPIR-V Runner fails to generate the serialized kernel argument data as the signature and constants representation has changed as part of recent update on main with PR - #3043 .

Prior to #3043 [Working Version]:

signature:
{0: '*fp32', 1: 'i32', 2: 'i32', 3: '*fp32', 4: 'i32', 5: 'i32', 6: '*fp32', 7: 'i32', 8: 'i32', 9: '*fp32', 10: 'i32', 11: 'i32'}

constants:
{2: 1, 5: 1, 8: 1, 11: 1, 12: 64, 13: 128, 14: 128, 15: False, 16: False, 17: False, 18: 'tf32', 19: False, 20: False, 21: False, 22: False, 23: triton.language.float32}

After #3043 [Failed version]:

Signature:
 {'X': '*fp32', 'stride_xm': 'i32', 'stride_xk': 'i32', 'Y': '*fp32', 'stride_yk': 'i32', 'stride_yn': 'i32', 'W': '*fp32', 'stride_wn': 'i32', 'stride_wl': 'i32', 'Z': '*fp32', 'stride_zm': 'i32', 'stride_zn': 'i32', 'BLOCK_M': 'constexpr', 'BLOCK_N': 'constexpr', 'BLOCK_K': 'constexpr', 'ADD_MATRIX': 'constexpr', 'ADD_ROWS': 'constexpr', 'ADD_COLS': 'constexpr', 'INPUT_PRECISION': 'constexpr', 'DO_SOFTMAX': 'constexpr', 'CHAIN_DOT': 'constexpr', 'COL_A': 'constexpr', 'COL_B': 'constexpr', 'out_dtype': 'constexpr'}

Constants:
{'BLOCK_M': 64, 'BLOCK_N': 128, 'BLOCK_K': 128, 'ADD_MATRIX': False, 'ADD_ROWS': False, 'ADD_COLS': False, 'INPUT_PRECISION': 'tf32', 'DO_SOFTMAX': False, 'CHAIN_DOT': False, 'COL_A': False, 'COL_B': False, 'out_dtype': triton.language.float32, 'stride_xk': 1, 'stride_yn': 1, 'stride_wl': 1, 'stride_zn': 1}

Environment details

Follow the SPIRV-Runner steps to reproduce this issue - https://github.com/intel/intel-xpu-backend-for-triton/tree/main/utils/SPIRVRunner

[kballeda]

@anmyachev / @whitneywhtsang : Just a general query on why we modified the way we structure signature and constants ?
In the snapshot attached here, the highlighted arguments are the valid constants for kernel.

[anmyachev]

@anmyachev / @whitneywhtsang : Just a general query on why we modified the way we structure signature and constants ? In the snapshot attached here, the highlighted arguments are the valid constants for kernel.

This is due to changes in the interface of Triton itself, in the following pull requests they were ported to our repository: #3043 and #3112.