Index: llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/nvptx-basic.ll =================================================================== --- /dev/null +++ llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/nvptx-basic.ll @@ -0,0 +1,42 @@ +; RUN: llc < %s -mtriple=nvptx-unknown-unknown | FileCheck %s + +%struct.St8x4 = type { [4 x i64] } + +define dso_local void @caller_St8x4(%struct.St8x4* nocapture noundef readonly byval(%struct.St8x4) align 8 %in, %struct.St8x4* nocapture noundef writeonly %ret) { + %call = tail call fastcc [4 x i64] @callee_St8x4(%struct.St8x4* noundef nonnull byval(%struct.St8x4) align 8 %in) + %.fca.0.extract = extractvalue [4 x i64] %call, 0 + %.fca.1.extract = extractvalue [4 x i64] %call, 1 + %.fca.2.extract = extractvalue [4 x i64] %call, 2 + %.fca.3.extract = extractvalue [4 x i64] %call, 3 + %ref.tmp.sroa.0.0..sroa_idx = getelementptr inbounds %struct.St8x4, %struct.St8x4* %ret, i64 0, i32 0, i64 0 + store i64 %.fca.0.extract, i64* %ref.tmp.sroa.0.0..sroa_idx, align 8 + %ref.tmp.sroa.4.0..sroa_idx3 = getelementptr inbounds %struct.St8x4, %struct.St8x4* %ret, i64 0, i32 0, i64 1 + store i64 %.fca.1.extract, i64* %ref.tmp.sroa.4.0..sroa_idx3, align 8 + %ref.tmp.sroa.5.0..sroa_idx5 = getelementptr inbounds %struct.St8x4, %struct.St8x4* %ret, i64 0, i32 0, i64 2 + store i64 %.fca.2.extract, i64* %ref.tmp.sroa.5.0..sroa_idx5, align 8 + %ref.tmp.sroa.6.0..sroa_idx7 = getelementptr inbounds %struct.St8x4, %struct.St8x4* %ret, i64 0, i32 0, i64 3 + store i64 %.fca.3.extract, i64* %ref.tmp.sroa.6.0..sroa_idx7, align 8 + ret void +} + + +define internal fastcc [4 x i64] @callee_St8x4(%struct.St8x4* nocapture noundef readonly byval(%struct.St8x4) align 8 %in) { + %arrayidx = getelementptr inbounds %struct.St8x4, %struct.St8x4* %in, i64 0, i32 0, i64 0 + %1 = load i64, i64* %arrayidx, align 8 + %arrayidx.1 = getelementptr inbounds %struct.St8x4, %struct.St8x4* %in, i64 0, i32 0, i64 1 + %2 = load i64, i64* %arrayidx.1, align 8 + %arrayidx.2 = getelementptr inbounds %struct.St8x4, %struct.St8x4* %in, i64 0, i32 0, i64 2 + %3 = load i64, i64* %arrayidx.2, align 8 + %arrayidx.3 = getelementptr inbounds %struct.St8x4, %struct.St8x4* %in, i64 0, i32 0, i64 3 + %4 = load i64, i64* %arrayidx.3, align 8 + %5 = insertvalue [4 x i64] poison, i64 %1, 0 + %6 = insertvalue [4 x i64] %5, i64 %2, 1 + %7 = insertvalue [4 x i64] %6, i64 %3, 2 + %oldret = insertvalue [4 x i64] %7, i64 %4, 3 + ret [4 x i64] %oldret +} + +define void @call_void() { + ret void +} + Index: llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/nvptx-basic.ll.expected =================================================================== --- /dev/null +++ llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/nvptx-basic.ll.expected @@ -0,0 +1,104 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py +; RUN: llc < %s -mtriple=nvptx-unknown-unknown | FileCheck %s + +%struct.St8x4 = type { [4 x i64] } + +define dso_local void @caller_St8x4(%struct.St8x4* nocapture noundef readonly byval(%struct.St8x4) align 8 %in, %struct.St8x4* nocapture noundef writeonly %ret) { +; CHECK-LABEL: caller_St8x4( +; CHECK: { +; CHECK-NEXT: .local .align 8 .b8 __local_depot0[32]; +; CHECK-NEXT: .reg .b32 %SP; +; CHECK-NEXT: .reg .b32 %SPL; +; CHECK-NEXT: .reg .b32 %r<4>; +; CHECK-NEXT: .reg .b64 %rd<17>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: mov.u32 %SPL, __local_depot0; +; CHECK-NEXT: cvta.local.u32 %SP, %SPL; +; CHECK-NEXT: ld.param.u32 %r1, [caller_St8x4_param_1]; +; CHECK-NEXT: add.u32 %r3, %SPL, 0; +; CHECK-NEXT: ld.param.u64 %rd1, [caller_St8x4_param_0+24]; +; CHECK-NEXT: ld.param.u64 %rd2, [caller_St8x4_param_0+16]; +; CHECK-NEXT: ld.param.u64 %rd3, [caller_St8x4_param_0+8]; +; CHECK-NEXT: ld.param.u64 %rd4, [caller_St8x4_param_0]; +; CHECK-NEXT: st.local.u64 [%r3], %rd4; +; CHECK-NEXT: st.local.u64 [%r3+8], %rd3; +; CHECK-NEXT: st.local.u64 [%r3+16], %rd2; +; CHECK-NEXT: st.local.u64 [%r3+24], %rd1; +; CHECK-NEXT: ld.u64 %rd5, [%SP+8]; +; CHECK-NEXT: ld.u64 %rd6, [%SP+0]; +; CHECK-NEXT: ld.u64 %rd7, [%SP+24]; +; CHECK-NEXT: ld.u64 %rd8, [%SP+16]; +; CHECK-NEXT: { // callseq 0, 0 +; CHECK-NEXT: .reg .b32 temp_param_reg; +; CHECK-NEXT: .param .align 16 .b8 param0[32]; +; CHECK-NEXT: st.param.v2.b64 [param0+0], {%rd6, %rd5}; +; CHECK-NEXT: st.param.v2.b64 [param0+16], {%rd8, %rd7}; +; CHECK-NEXT: .param .align 16 .b8 retval0[32]; +; CHECK-NEXT: call.uni (retval0), +; CHECK-NEXT: callee_St8x4, +; CHECK-NEXT: ( +; CHECK-NEXT: param0 +; CHECK-NEXT: ); +; CHECK-NEXT: ld.param.v2.b64 {%rd9, %rd10}, [retval0+0]; +; CHECK-NEXT: ld.param.v2.b64 {%rd11, %rd12}, [retval0+16]; +; CHECK-NEXT: } // callseq 0 +; CHECK-NEXT: st.u64 [%r1], %rd9; +; CHECK-NEXT: st.u64 [%r1+8], %rd10; +; CHECK-NEXT: st.u64 [%r1+16], %rd11; +; CHECK-NEXT: st.u64 [%r1+24], %rd12; +; CHECK-NEXT: ret; + %call = tail call fastcc [4 x i64] @callee_St8x4(%struct.St8x4* noundef nonnull byval(%struct.St8x4) align 8 %in) + %.fca.0.extract = extractvalue [4 x i64] %call, 0 + %.fca.1.extract = extractvalue [4 x i64] %call, 1 + %.fca.2.extract = extractvalue [4 x i64] %call, 2 + %.fca.3.extract = extractvalue [4 x i64] %call, 3 + %ref.tmp.sroa.0.0..sroa_idx = getelementptr inbounds %struct.St8x4, %struct.St8x4* %ret, i64 0, i32 0, i64 0 + store i64 %.fca.0.extract, i64* %ref.tmp.sroa.0.0..sroa_idx, align 8 + %ref.tmp.sroa.4.0..sroa_idx3 = getelementptr inbounds %struct.St8x4, %struct.St8x4* %ret, i64 0, i32 0, i64 1 + store i64 %.fca.1.extract, i64* %ref.tmp.sroa.4.0..sroa_idx3, align 8 + %ref.tmp.sroa.5.0..sroa_idx5 = getelementptr inbounds %struct.St8x4, %struct.St8x4* %ret, i64 0, i32 0, i64 2 + store i64 %.fca.2.extract, i64* %ref.tmp.sroa.5.0..sroa_idx5, align 8 + %ref.tmp.sroa.6.0..sroa_idx7 = getelementptr inbounds %struct.St8x4, %struct.St8x4* %ret, i64 0, i32 0, i64 3 + store i64 %.fca.3.extract, i64* %ref.tmp.sroa.6.0..sroa_idx7, align 8 + ret void +} + + +define internal fastcc [4 x i64] @callee_St8x4(%struct.St8x4* nocapture noundef readonly byval(%struct.St8x4) align 8 %in) { +; CHECK-LABEL: callee_St8x4( +; CHECK: // @callee_St8x4 +; CHECK-NEXT: { +; CHECK-NEXT: .reg .b64 %rd<5>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.v2.u64 {%rd1, %rd2}, [callee_St8x4_param_0]; +; CHECK-NEXT: ld.param.v2.u64 {%rd3, %rd4}, [callee_St8x4_param_0+16]; +; CHECK-NEXT: st.param.v2.b64 [func_retval0+0], {%rd1, %rd2}; +; CHECK-NEXT: st.param.v2.b64 [func_retval0+16], {%rd3, %rd4}; +; CHECK-NEXT: ret; + %arrayidx = getelementptr inbounds %struct.St8x4, %struct.St8x4* %in, i64 0, i32 0, i64 0 + %1 = load i64, i64* %arrayidx, align 8 + %arrayidx.1 = getelementptr inbounds %struct.St8x4, %struct.St8x4* %in, i64 0, i32 0, i64 1 + %2 = load i64, i64* %arrayidx.1, align 8 + %arrayidx.2 = getelementptr inbounds %struct.St8x4, %struct.St8x4* %in, i64 0, i32 0, i64 2 + %3 = load i64, i64* %arrayidx.2, align 8 + %arrayidx.3 = getelementptr inbounds %struct.St8x4, %struct.St8x4* %in, i64 0, i32 0, i64 3 + %4 = load i64, i64* %arrayidx.3, align 8 + %5 = insertvalue [4 x i64] poison, i64 %1, 0 + %6 = insertvalue [4 x i64] %5, i64 %2, 1 + %7 = insertvalue [4 x i64] %6, i64 %3, 2 + %oldret = insertvalue [4 x i64] %7, i64 %4, 3 + ret [4 x i64] %oldret +} + +define void @call_void() { +; CHECK-LABEL: call_void( +; CHECK: { +; CHECK-EMPTY: +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ret; + ret void +} + Index: llvm/test/tools/UpdateTestChecks/update_llc_test_checks/nvptx-basic.test =================================================================== --- /dev/null +++ llvm/test/tools/UpdateTestChecks/update_llc_test_checks/nvptx-basic.test @@ -0,0 +1,4 @@ +# REQUIRES: nvptx-registered-target + +# RUN: cp -f %S/Inputs/nvptx-basic.ll %t.ll && %update_llc_test_checks %t.ll +# RUN: diff -u %S/Inputs/nvptx-basic.ll.expected %t.ll Index: llvm/utils/UpdateTestChecks/asm.py =================================================================== --- llvm/utils/UpdateTestChecks/asm.py +++ llvm/utils/UpdateTestChecks/asm.py @@ -178,6 +178,27 @@ r'.Lfunc_end[0-9]+:\n', flags=(re.M | re.S)) +ASM_FUNCTION_NVPTX_RE = re.compile( + # function attributes and retval + # .visible .func (.param .align 16 .b8 func_retval0[32]) + r'^(\.visible\s+)?\.func\s+(\([^\)]*\)\s*)?' + + # function name + r'(?P[^\(\n]+)' + + # function parameters + # ( + # .param .align 16 .b8 callee_St8x4_param_0[32] + # ) // -- Begin function callee_St8x4 + r'\([^\)]*\)(\s*//[^\n]*)?\n' + + # function body + r'(?P.*?)\n' + + # function body end marker + r'\s*// -- End function', + flags=(re.M | re.S)) + SCRUB_X86_SHUFFLES_RE = ( re.compile( r'^(\s*\w+) [^#\n]+#+ ((?:[xyz]mm\d+|mem)( \{%k\d+\}( \{z\})?)? = .*)$', @@ -388,44 +409,55 @@ asm = common.SCRUB_TRAILING_WHITESPACE_RE.sub(r'', asm) return asm +def scrub_asm_nvptx(asm, args): + # Scrub runs of whitespace out of the assembly, but leave the leading + # whitespace in place. + asm = common.SCRUB_WHITESPACE_RE.sub(r' ', asm) + # Expand the tabs used for indentation. + asm = string.expandtabs(asm, 2) + # Strip trailing whitespace. + asm = common.SCRUB_TRAILING_WHITESPACE_RE.sub(r'', asm) + return asm + def get_run_handler(triple): target_handlers = { - 'i686': (scrub_asm_x86, ASM_FUNCTION_X86_RE), - 'x86': (scrub_asm_x86, ASM_FUNCTION_X86_RE), - 'i386': (scrub_asm_x86, ASM_FUNCTION_X86_RE), - 'arm64_32-apple-ios': (scrub_asm_arm_eabi, ASM_FUNCTION_AARCH64_DARWIN_RE), - 'aarch64': (scrub_asm_arm_eabi, ASM_FUNCTION_AARCH64_RE), - 'aarch64-apple-darwin': (scrub_asm_arm_eabi, ASM_FUNCTION_AARCH64_DARWIN_RE), - 'aarch64-apple-ios': (scrub_asm_arm_eabi, ASM_FUNCTION_AARCH64_DARWIN_RE), - 'hexagon': (scrub_asm_hexagon, ASM_FUNCTION_HEXAGON_RE), - 'r600': (scrub_asm_amdgpu, ASM_FUNCTION_AMDGPU_RE), - 'amdgcn': (scrub_asm_amdgpu, ASM_FUNCTION_AMDGPU_RE), - 'arm': (scrub_asm_arm_eabi, ASM_FUNCTION_ARM_RE), - 'arm64': (scrub_asm_arm_eabi, ASM_FUNCTION_AARCH64_RE), - 'arm64e': (scrub_asm_arm_eabi, ASM_FUNCTION_AARCH64_DARWIN_RE), - 'arm64-apple-ios': (scrub_asm_arm_eabi, ASM_FUNCTION_AARCH64_DARWIN_RE), - 'armv7-apple-ios' : (scrub_asm_arm_eabi, ASM_FUNCTION_ARM_IOS_RE), - 'armv7-apple-darwin': (scrub_asm_arm_eabi, ASM_FUNCTION_ARM_DARWIN_RE), - 'thumb': (scrub_asm_arm_eabi, ASM_FUNCTION_ARM_RE), - 'thumb-macho': (scrub_asm_arm_eabi, ASM_FUNCTION_ARM_MACHO_RE), - 'thumbv5-macho': (scrub_asm_arm_eabi, ASM_FUNCTION_ARM_MACHO_RE), - 'thumbv7s-apple-darwin' : (scrub_asm_arm_eabi, ASM_FUNCTION_THUMBS_DARWIN_RE), - 'thumbv7-apple-darwin' : (scrub_asm_arm_eabi, ASM_FUNCTION_THUMB_DARWIN_RE), - 'thumbv7-apple-ios' : (scrub_asm_arm_eabi, ASM_FUNCTION_ARM_IOS_RE), - 'm68k': (scrub_asm_m68k, ASM_FUNCTION_M68K_RE), - 'mips': (scrub_asm_mips, ASM_FUNCTION_MIPS_RE), - 'msp430': (scrub_asm_msp430, ASM_FUNCTION_MSP430_RE), - 'avr': (scrub_asm_avr, ASM_FUNCTION_AVR_RE), - 'ppc32': (scrub_asm_powerpc, ASM_FUNCTION_PPC_RE), - 'powerpc': (scrub_asm_powerpc, ASM_FUNCTION_PPC_RE), - 'riscv32': (scrub_asm_riscv, ASM_FUNCTION_RISCV_RE), - 'riscv64': (scrub_asm_riscv, ASM_FUNCTION_RISCV_RE), - 'lanai': (scrub_asm_lanai, ASM_FUNCTION_LANAI_RE), - 'sparc': (scrub_asm_sparc, ASM_FUNCTION_SPARC_RE), - 's390x': (scrub_asm_systemz, ASM_FUNCTION_SYSTEMZ_RE), - 'wasm32': (scrub_asm_wasm32, ASM_FUNCTION_WASM32_RE), - 've': (scrub_asm_ve, ASM_FUNCTION_VE_RE), - 'csky': (scrub_asm_csky, ASM_FUNCTION_CSKY_RE), + 'i686': (scrub_asm_x86, ASM_FUNCTION_X86_RE, ':'), + 'x86': (scrub_asm_x86, ASM_FUNCTION_X86_RE, ':'), + 'i386': (scrub_asm_x86, ASM_FUNCTION_X86_RE, ':'), + 'arm64_32-apple-ios': (scrub_asm_arm_eabi, ASM_FUNCTION_AARCH64_DARWIN_RE, ':'), + 'aarch64': (scrub_asm_arm_eabi, ASM_FUNCTION_AARCH64_RE, ':'), + 'aarch64-apple-darwin': (scrub_asm_arm_eabi, ASM_FUNCTION_AARCH64_DARWIN_RE, ':'), + 'aarch64-apple-ios': (scrub_asm_arm_eabi, ASM_FUNCTION_AARCH64_DARWIN_RE, ':'), + 'hexagon': (scrub_asm_hexagon, ASM_FUNCTION_HEXAGON_RE, ':'), + 'r600': (scrub_asm_amdgpu, ASM_FUNCTION_AMDGPU_RE, ':'), + 'amdgcn': (scrub_asm_amdgpu, ASM_FUNCTION_AMDGPU_RE, ':'), + 'arm': (scrub_asm_arm_eabi, ASM_FUNCTION_ARM_RE, ':'), + 'arm64': (scrub_asm_arm_eabi, ASM_FUNCTION_AARCH64_RE, ':'), + 'arm64e': (scrub_asm_arm_eabi, ASM_FUNCTION_AARCH64_DARWIN_RE, ':'), + 'arm64-apple-ios': (scrub_asm_arm_eabi, ASM_FUNCTION_AARCH64_DARWIN_RE, ':'), + 'armv7-apple-ios' : (scrub_asm_arm_eabi, ASM_FUNCTION_ARM_IOS_RE, ':'), + 'armv7-apple-darwin': (scrub_asm_arm_eabi, ASM_FUNCTION_ARM_DARWIN_RE, ':'), + 'thumb': (scrub_asm_arm_eabi, ASM_FUNCTION_ARM_RE, ':'), + 'thumb-macho': (scrub_asm_arm_eabi, ASM_FUNCTION_ARM_MACHO_RE, ':'), + 'thumbv5-macho': (scrub_asm_arm_eabi, ASM_FUNCTION_ARM_MACHO_RE, ':'), + 'thumbv7s-apple-darwin' : (scrub_asm_arm_eabi, ASM_FUNCTION_THUMBS_DARWIN_RE, ':'), + 'thumbv7-apple-darwin' : (scrub_asm_arm_eabi, ASM_FUNCTION_THUMB_DARWIN_RE, ':'), + 'thumbv7-apple-ios' : (scrub_asm_arm_eabi, ASM_FUNCTION_ARM_IOS_RE, ':'), + 'm68k': (scrub_asm_m68k, ASM_FUNCTION_M68K_RE, ':'), + 'mips': (scrub_asm_mips, ASM_FUNCTION_MIPS_RE, ':'), + 'msp430': (scrub_asm_msp430, ASM_FUNCTION_MSP430_RE, ':'), + 'avr': (scrub_asm_avr, ASM_FUNCTION_AVR_RE, ':'), + 'ppc32': (scrub_asm_powerpc, ASM_FUNCTION_PPC_RE, ':'), + 'powerpc': (scrub_asm_powerpc, ASM_FUNCTION_PPC_RE, ':'), + 'riscv32': (scrub_asm_riscv, ASM_FUNCTION_RISCV_RE, ':'), + 'riscv64': (scrub_asm_riscv, ASM_FUNCTION_RISCV_RE, ':'), + 'lanai': (scrub_asm_lanai, ASM_FUNCTION_LANAI_RE, ':'), + 'sparc': (scrub_asm_sparc, ASM_FUNCTION_SPARC_RE, ':'), + 's390x': (scrub_asm_systemz, ASM_FUNCTION_SYSTEMZ_RE, ':'), + 'wasm32': (scrub_asm_wasm32, ASM_FUNCTION_WASM32_RE, ':'), + 've': (scrub_asm_ve, ASM_FUNCTION_VE_RE, ':'), + 'csky': (scrub_asm_csky, ASM_FUNCTION_CSKY_RE, ':'), + 'nvptx': (scrub_asm_nvptx, ASM_FUNCTION_NVPTX_RE, '(') } handler = None best_prefix = '' @@ -442,10 +474,11 @@ ##### Generator of assembly CHECK lines def add_checks(output_lines, comment_marker, prefix_list, func_dict, - func_name, is_filtered): + func_name, is_filtered, label_suffix_dict): # Label format is based on ASM string. - check_label_format = '{} %s-LABEL: %s%s:'.format(comment_marker) + check_label_format = '{} %s-LABEL: %s%s%s'.format(comment_marker) global_vars_seen_dict = {} common.add_checks(output_lines, comment_marker, prefix_list, func_dict, func_name, check_label_format, True, False, - global_vars_seen_dict, is_filtered = is_filtered) + global_vars_seen_dict, is_filtered=is_filtered, + label_suffix_dict=label_suffix_dict) Index: llvm/utils/UpdateTestChecks/common.py =================================================================== --- llvm/utils/UpdateTestChecks/common.py +++ llvm/utils/UpdateTestChecks/common.py @@ -455,6 +455,7 @@ # Strip double-quotes if input was read by UTC_ARGS self._replace_value_regex = list(map(lambda x: x.strip('"'), flags.replace_value_regex)) self._func_dict = {} + self._label_suffix_dict = {} self._func_order = {} self._global_var_dict = {} for tuple in run_list: @@ -468,6 +469,9 @@ warn('Prefix %s had conflicting output from different RUN lines for all functions in test %s' % (prefix,self._path,)) return self._func_dict + def label_suffix_dict(self): + return self._label_suffix_dict + def func_order(self): return self._func_order @@ -477,7 +481,7 @@ def is_filtered(self): return bool(self._filters) - def process_run_line(self, function_re, scrubber, raw_tool_output, prefixes, is_backend): + def process_run_line(self, function_re, scrubber, raw_tool_output, prefixes, is_backend, label_suffix=None): build_global_values_dictionary(self._global_var_dict, raw_tool_output, prefixes) for m in function_re.finditer(raw_tool_output): if not m: @@ -559,6 +563,7 @@ self._func_dict[prefix][func] = function_body( scrubbed_body, scrubbed_extra, args_and_sig, attrs) + self._label_suffix_dict[prefix] = label_suffix self._func_order[prefix].append(func) def _get_failed_prefixes(self): @@ -783,7 +788,7 @@ return lines -def add_checks(output_lines, comment_marker, prefix_list, func_dict, func_name, check_label_format, is_backend, is_analyze, global_vars_seen_dict, is_filtered): +def add_checks(output_lines, comment_marker, prefix_list, func_dict, func_name, check_label_format, is_backend, is_analyze, global_vars_seen_dict, is_filtered, label_suffix_dict=None): # prefix_exclusions are prefixes we cannot use to print the function because it doesn't exist in run lines that use these prefixes as well. prefix_exclusions = set() printed_prefixes = [] @@ -836,10 +841,16 @@ args_and_sig = str(func_dict[checkprefix][func_name].args_and_sig) args_and_sig = generalize_check_lines([args_and_sig], is_analyze, vars_seen, global_vars_seen)[0] if '[[' in args_and_sig: - output_lines.append(check_label_format % (checkprefix, func_name, '')) + if label_suffix_dict is not None: + output_lines.append(check_label_format % (checkprefix, func_name, '', label_suffix_dict[checkprefix])) + else: + output_lines.append(check_label_format % (checkprefix, func_name, '')) output_lines.append('%s %s-SAME: %s' % (comment_marker, checkprefix, args_and_sig)) else: - output_lines.append(check_label_format % (checkprefix, func_name, args_and_sig)) + if label_suffix_dict is not None: + output_lines.append(check_label_format % (checkprefix, func_name, args_and_sig, label_suffix_dict[checkprefix])) + else: + output_lines.append(check_label_format % (checkprefix, func_name, args_and_sig)) func_body = str(func_dict[checkprefix][func_name]).splitlines() if not func_body: # We have filtered everything. Index: llvm/utils/UpdateTestChecks/isel.py =================================================================== --- llvm/utils/UpdateTestChecks/isel.py +++ llvm/utils/UpdateTestChecks/isel.py @@ -42,16 +42,18 @@ if handler is None: common.debug('Using default handler.') - handler = (scrub_isel_default, ISEL_FUNCTION_DEFAULT_RE) + handler = (scrub_isel_default, ISEL_FUNCTION_DEFAULT_RE, ':') return handler ##### Generator of iSel CHECK lines -def add_checks(output_lines, comment_marker, prefix_list, func_dict, func_name, is_filtered): +def add_checks(output_lines, comment_marker, prefix_list, func_dict, + func_name, is_filtered, label_suffix_dict): # Label format is based on iSel string. - check_label_format = '{} %s-LABEL: %s%s:'.format(comment_marker) + check_label_format = '{} %s-LABEL: %s%s%s'.format(comment_marker) global_vars_seen_dict = {} common.add_checks(output_lines, comment_marker, prefix_list, func_dict, func_name, check_label_format, True, False, - global_vars_seen_dict, is_filtered = is_filtered) + global_vars_seen_dict, is_filtered=is_filtered, + label_suffix_dict=label_suffix_dict) Index: llvm/utils/update_llc_test_checks.py =================================================================== --- llvm/utils/update_llc_test_checks.py +++ llvm/utils/update_llc_test_checks.py @@ -135,8 +135,8 @@ if not triple: triple = common.get_triple_from_march(march_in_cmd) - scrubber, function_re = output_type.get_run_handler(triple) - builder.process_run_line(function_re, scrubber, raw_tool_output, prefixes, True) + scrubber, function_re, label_suffix = output_type.get_run_handler(triple) + builder.process_run_line(function_re, scrubber, raw_tool_output, prefixes, True, label_suffix) func_dict = builder.finish_and_get_func_dict() @@ -167,9 +167,10 @@ check_indent + ';', lambda my_output_lines, prefixes, func: output_type.add_checks(my_output_lines, - check_indent + ';', - prefixes, func_dict, func, - is_filtered=builder.is_filtered())) + check_indent + ';', + prefixes, func_dict, func, + is_filtered=builder.is_filtered(), + label_suffix_dict=builder.label_suffix_dict())) else: for input_info in ti.iterlines(output_lines): input_line = input_info.line @@ -185,8 +186,9 @@ # Print out the various check lines here. output_type.add_checks(output_lines, check_indent + ';', run_list, - func_dict, func_name, - is_filtered=builder.is_filtered()) + func_dict, func_name, + is_filtered=builder.is_filtered(), + label_suffix_dict=builder.label_suffix_dict()) is_in_function_start = False if is_in_function: