[UpdateTestChecks] Add NVPTX support in update_llc_test_checks.py

[kovdan01]

Summary

This patch makes possible generating NVPTX assembly check lines with
update_llc_test_checks.py utility.

[tra]

llvm/utils/UpdateTestChecks/asm.py

184

There may be other attributes (e.g. .weak) and they may come in different order.
It may be better to use something like (\.(func|visible|weak|entry|noreturn|extern)\s+).

430

It would be great to document what the function is expected to return.

503

The third field is a label_suffix, but it's not clear why it's ( for NVPTX.
AFAICT, PTX does use ':' for the labels: https://cuda.godbolt.org/z/34185xEbY

Looks like it's supposed to serve as a terminator for the CHECK-LABEL (i.e. it's a suffix in terms of FileCheck, but not necessarily in terms of the actual asm labels).
This may be something to add to the description of the function return value.

I wonder if we can incorporate this suffix into a named regex match group, so we don't have to pass it explicitly.

[kovdan01]

llvm/utils/UpdateTestChecks/asm.py

503

The third field is a label_suffix, but it's not clear why it's ( for NVPTX.
AFAICT, PTX does use ':' for the labels: https://cuda.godbolt.org/z/34185xEbY

I suppose that I used a wrong term for naming that, and maybe I should rename that to function_decl_suffix or something like this.
Maybe I have missed something, but as far as I understood functions in PTX begin like this - having ( after function name:

.visible .entry square(int *, int)(
        .param .u64 square(int *, int)_param_0,
        .param .u32 square(int *, int)_param_1
)

In other assemblies, function "declarations" are just labels looking like foo:. So I introduced this new suffix field, because we need to have some marker of function start. Previously, a colon after label was hardcoded.

I wonder if we can incorporate this suffix into a named regex match group, so we don't have to pass it explicitly.

Unfortunately, that seems to be even more non-transparent - this was the first way that I tried and failed. Regex named group func is intended to hold the same function name as IR function has. Here is code from add_checks function from common.py:

# If we do not have output for this prefix we skip it.
if not func_dict[checkprefix][func_name]:
  continue

I also don't like that separate suffix field, but have no idea how to get rid of it in a nice way.

[tra]

llvm/utils/UpdateTestChecks/asm.py

503

I wonder if we can incorporate this suffix into a named regex match group, so we don't have to pass it explicitly.

Unfortunately, that seems to be even more non-transparent - this was the first way that I tried and failed. Regex named group func is intended to hold the same function name as IR function has. Here is code from add_checks function from common.py:

# If we do not have output for this prefix we skip it.
if not func_dict[checkprefix][func_name]:
  continue

I was thinking of adding the separator as yet another named group within the existing RE for function parameters:
r'\([^\)]*\)(\s*//[^\n]*)?\n' -> r'(?P<separator>\()[^\)]*\)(\s*//[^\n]*)?\n'

The named group 'func' remains unchanged, but now you can extract the separator via its own named group, save it along with the function name and later retrieve and use it in add_checks.

To minimize churn, the separator group may be optional and would default to ':' which would work for existing targets.

[tra]

LGTM, but I'll defer to UpdateTestChecks owners for the final approval.

llvm/utils/UpdateTestChecks/common.py

493–497

Existence of the named group can be checked directly:

func_name_separator = m.group('func_name_separator')  if 'func_name_separator' in m.groupdict() else ":"

[kovdan01]

llvm/utils/UpdateTestChecks/common.py

493–497

Thanks for useful advice! Fixed.

[kovdan01]

@lattner Hello Chris, could you please suggest a reviewer for this patch? I'm not sure who is the owner of this part of repository. Thanks!

[lattner]

I don't know who would be the best to review this, @MaskRay do you know?

[MaskRay]

llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/nvptx-basic.ll

21

nit: keep just one blank line

42

nit: delete trailing blank line

llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/nvptx-basic.ll.expected

81

opaque pointers have been enabled today.
You may need to switch to ptr

[kovdan01]

llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/nvptx-basic.ll.expected

81

Fixed, thanks!

llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/nvptx-basic.ll

+; 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
+}
+

llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/nvptx-basic.ll.expected

+; 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
+}
+

llvm/test/tools/UpdateTestChecks/update_llc_test_checks/nvptx-basic.test

+# 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

llvm/utils/UpdateTestChecks/asm.py

ASM_FUNCTION_VE_RE = re.compile(
     flags=(re.M | re.S))
 
 ASM_FUNCTION_CSKY_RE = re.compile(
     r'^_?(?P<func>[^:]+):[ \t]*#+[ \t]*@(?P=func)\n(?:\s*\.?Lfunc_begin[^:\n]*:\n)?[^:]*?'
     r'(?P<body>^##?[ \t]+[^:]+:.*?)\s*'
     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*)?'
+    r'^(\.(func|visible|weak|entry|noreturn|extern)\s+)+(\([^\)]*\)\s*)?'
+
+    # function name
+    r'(?P<func>[^\(\n]+)'
+
+    # function name separator (opening brace)
+    r'(?P<func_name_separator>\()'
+
+    # function parameters
+    # (
+    #   .param .align 16 .b8 callee_St8x4_param_0[32]
+    # ) // -- Begin function callee_St8x4
+    r'[^\)]*\)(\s*//[^\n]*)?\n'
+
+    # function body
+    r'(?P<body>.*?)\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\})?)? = .*)$',
         flags=re.M))
 
 SCRUB_X86_SHUFFLES_NO_MEM_RE = (
     re.compile(
         r'^(\s*\w+) [^#\n]+#+ ((?:[xyz]mm\d+|mem)( \{%k\d+\}( \{z\})?)? = (?!.*(?:mem)).*)$',
def scrub_asm_csky(asm, args):
   # Expand the tabs used for indentation.
   asm = string.expandtabs(asm, 2)
   # Strip kill operands inserted into the asm.
   asm = common.SCRUB_KILL_COMMENT_RE.sub('', asm)
   # Strip trailing whitespace.
   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
+
+
+# Returns a tuple of a scrub function and a function regex. Scrub function is
+# used to alter function body in some way, for example, remove traling spaces.
+# Function regex is used to match function name, body, etc. in raw llc output.
 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),
target_handlers = {
       '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 = ''
   for prefix, s in target_handlers.items():
     if triple.startswith(prefix) and len(prefix) > len(best_prefix):
       handler = s
       best_prefix = prefix
 
   if handler is None:
     raise KeyError('Triple %r is not supported' % (triple))
 
   return handler
 
 ##### Generator of assembly CHECK lines
 
 def add_checks(output_lines, comment_marker, prefix_list, func_dict,
                    func_name, is_filtered):
   # 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)

llvm/utils/UpdateTestChecks/common.py

def do_scrub(body, scrubber, scrubber_args, extra):
   if scrubber_args:
     local_args = copy.deepcopy(scrubber_args)
     local_args[0].extra_scrub = extra
     return scrubber(body, *local_args)
   return scrubber(body, *scrubber_args)
 
 # Build up a dictionary of all the function bodies.
 class function_body(object):
-  def __init__(self, string, extra, args_and_sig, attrs):
+  def __init__(self, string, extra, args_and_sig, attrs, func_name_separator):
     self.scrub = string
     self.extrascrub = extra
     self.args_and_sig = args_and_sig
     self.attrs = attrs
+    self.func_name_separator = func_name_separator
   def is_same_except_arg_names(self, extrascrub, args_and_sig, attrs, is_backend):
     arg_names = set()
     def drop_arg_names(match):
         arg_names.add(match.group(variable_group_in_ir_value_match))
         if match.group(attribute_group_in_ir_value_match):
             attr = match.group(attribute_group_in_ir_value_match)
         else:
             attr = ''
class FunctionTestBuilder:
 
   def process_run_line(self, function_re, scrubber, raw_tool_output, prefixes, is_backend):
     build_global_values_dictionary(self._global_var_dict, raw_tool_output, prefixes)
     for m in function_re.finditer(raw_tool_output):
       if not m:
         continue
       func = m.group('func')
       body = m.group('body')
+      # func_name_separator is the string that is placed right after function name at the
+      # beginning of assembly function definition. In most assemblies, that is just a
+      # colon: `foo:`. But, for example, in nvptx it is a brace: `foo(`. If is_backend is
+      # False, just assume that separator is an empty string.
+      if is_backend:
+        # Use ':' as default separator.
+        func_name_separator = m.group('func_name_separator') if 'func_name_separator' in m.groupdict() else ':'
+      else:
+        func_name_separator = ''
       attrs = m.group('attrs') if self._check_attributes else ''
       # Determine if we print arguments, the opening brace, or nothing after the
       # function name
       if self._record_args and 'args_and_sig' in m.groupdict():
           args_and_sig = scrub_body(m.group('args_and_sig').strip())
       elif 'args_and_sig' in m.groupdict():
           args_and_sig = '('
       else:
           args_and_sig = ''
for m in function_re.finditer(raw_tool_output):
               # This means a previous RUN line produced a body for this function
               # that is different from the one produced by this current RUN line,
               # so the body can't be common accross RUN lines. We use None to
               # indicate that.
               self._func_dict[prefix][func] = None
               continue
 
         self._func_dict[prefix][func] = function_body(
-            scrubbed_body, scrubbed_extra, args_and_sig, attrs)
+            scrubbed_body, scrubbed_extra, args_and_sig, attrs, func_name_separator)
         self._func_order[prefix].append(func)
 
   def _get_failed_prefixes(self):
     # This returns the list of those prefixes that failed to match any function,
     # because there were conflicting bodies produced by different RUN lines, in
     # all instances of the prefix. Effectively, this prefix is unused and should
     # be removed.
     for prefix in self._func_dict:
for checkprefix in checkprefixes:
       vars_seen = set()
       printed_prefixes.append(checkprefix)
       attrs = str(func_dict[checkprefix][func_name].attrs)
       attrs = '' if attrs == 'None' else attrs
       if attrs:
         output_lines.append('%s %s: Function Attrs: %s' % (comment_marker, checkprefix, attrs))
       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]
+      func_name_separator = func_dict[checkprefix][func_name].func_name_separator
       if '[[' in args_and_sig:
-        output_lines.append(check_label_format % (checkprefix, func_name, ''))
+        output_lines.append(check_label_format % (checkprefix, func_name, '', func_name_separator))
         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))
+        output_lines.append(check_label_format % (checkprefix, func_name, args_and_sig, func_name_separator))
       func_body = str(func_dict[checkprefix][func_name]).splitlines()
       if not func_body:
         # We have filtered everything.
         continue
 
       # For ASM output, just emit the check lines.
       if is_backend:
         body_start = 1
for checkprefix in checkprefixes:
               global_vars_seen_dict[checkprefix][key] = global_vars_seen[key]
       break
 
 def add_ir_checks(output_lines, comment_marker, prefix_list, func_dict,
                   func_name, preserve_names, function_sig,
                   global_vars_seen_dict, is_filtered):
   # Label format is based on IR string.
   function_def_regex = 'define {{[^@]+}}' if function_sig else ''
-  check_label_format = '{} %s-LABEL: {}@%s%s'.format(comment_marker, function_def_regex)
+  check_label_format = '{} %s-LABEL: {}@%s%s%s'.format(comment_marker, function_def_regex)
   add_checks(output_lines, comment_marker, prefix_list, func_dict, func_name,
              check_label_format, False, preserve_names, global_vars_seen_dict,
              is_filtered)
 
 def add_analyze_checks(output_lines, comment_marker, prefix_list, func_dict, func_name, is_filtered):
-  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 = {}
   add_checks(output_lines, comment_marker, prefix_list, func_dict, func_name,
              check_label_format, False, True, global_vars_seen_dict,
              is_filtered)
 
 def build_global_values_dictionary(glob_val_dict, raw_tool_output, prefixes):
   for nameless_value in nameless_values:
     if nameless_value.global_ir_prefix is None:

llvm/utils/UpdateTestChecks/isel.py

if handler is None:
     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):
   # 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)