[llvm] 8e43cba - [UpdateTestChecks] Add NVPTX support in update_llc_test_checks.py
Daniil Kovalev via llvm-commits
llvm-commits at lists.llvm.org
Fri Apr 15 01:02:32 PDT 2022
Author: Daniil Kovalev
Date: 2022-04-15T11:01:53+03:00
New Revision: 8e43cbab33765c476337571e5ed11b005199dd0d
URL: https://github.com/llvm/llvm-project/commit/8e43cbab33765c476337571e5ed11b005199dd0d
DIFF: https://github.com/llvm/llvm-project/commit/8e43cbab33765c476337571e5ed11b005199dd0d.diff
LOG: [UpdateTestChecks] Add NVPTX support in update_llc_test_checks.py
This patch makes possible generating NVPTX assembly check lines with
update_llc_test_checks.py utility.
Differential Revision: https://reviews.llvm.org/D122986
Added:
llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/nvptx-basic.ll
llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/nvptx-basic.ll.expected
llvm/test/tools/UpdateTestChecks/update_llc_test_checks/nvptx-basic.test
Modified:
llvm/utils/UpdateTestChecks/asm.py
llvm/utils/UpdateTestChecks/common.py
llvm/utils/UpdateTestChecks/isel.py
Removed:
################################################################################
diff --git a/llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/nvptx-basic.ll b/llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/nvptx-basic.ll
new file mode 100644
index 0000000000000..31fe9e96db30f
--- /dev/null
+++ b/llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/nvptx-basic.ll
@@ -0,0 +1,38 @@
+; RUN: llc < %s -mtriple=nvptx-unknown-unknown | FileCheck %s
+
+%struct.St8x4 = type { [4 x i64] }
+
+define dso_local void @caller_St8x4(ptr nocapture noundef readonly byval(%struct.St8x4) align 8 %in, ptr nocapture noundef writeonly %ret) {
+ %call = tail call fastcc [4 x i64] @callee_St8x4(ptr noundef nonnull byval(%struct.St8x4) align 8 %in) #2
+ %.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
+ store i64 %.fca.0.extract, ptr %ret, align 8
+ %ref.tmp.sroa.4.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 8
+ store i64 %.fca.1.extract, ptr %ref.tmp.sroa.4.0..sroa_idx, align 8
+ %ref.tmp.sroa.5.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 16
+ store i64 %.fca.2.extract, ptr %ref.tmp.sroa.5.0..sroa_idx, align 8
+ %ref.tmp.sroa.6.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 24
+ store i64 %.fca.3.extract, ptr %ref.tmp.sroa.6.0..sroa_idx, align 8
+ ret void
+}
+
+define internal fastcc [4 x i64] @callee_St8x4(ptr nocapture noundef readonly byval(%struct.St8x4) align 8 %in) {
+ %1 = load i64, ptr %in, align 8
+ %arrayidx.1 = getelementptr inbounds [4 x i64], ptr %in, i64 0, i64 1
+ %2 = load i64, ptr %arrayidx.1, align 8
+ %arrayidx.2 = getelementptr inbounds [4 x i64], ptr %in, i64 0, i64 2
+ %3 = load i64, ptr %arrayidx.2, align 8
+ %arrayidx.3 = getelementptr inbounds [4 x i64], ptr %in, i64 0, i64 3
+ %4 = load i64, ptr %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
+}
diff --git a/llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/nvptx-basic.ll.expected b/llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/nvptx-basic.ll.expected
new file mode 100644
index 0000000000000..a65f140b46015
--- /dev/null
+++ b/llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/nvptx-basic.ll.expected
@@ -0,0 +1,100 @@
+; 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(ptr nocapture noundef readonly byval(%struct.St8x4) align 8 %in, ptr 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(ptr noundef nonnull byval(%struct.St8x4) align 8 %in) #2
+ %.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
+ store i64 %.fca.0.extract, ptr %ret, align 8
+ %ref.tmp.sroa.4.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 8
+ store i64 %.fca.1.extract, ptr %ref.tmp.sroa.4.0..sroa_idx, align 8
+ %ref.tmp.sroa.5.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 16
+ store i64 %.fca.2.extract, ptr %ref.tmp.sroa.5.0..sroa_idx, align 8
+ %ref.tmp.sroa.6.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 24
+ store i64 %.fca.3.extract, ptr %ref.tmp.sroa.6.0..sroa_idx, align 8
+ ret void
+}
+
+define internal fastcc [4 x i64] @callee_St8x4(ptr 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;
+ %1 = load i64, ptr %in, align 8
+ %arrayidx.1 = getelementptr inbounds [4 x i64], ptr %in, i64 0, i64 1
+ %2 = load i64, ptr %arrayidx.1, align 8
+ %arrayidx.2 = getelementptr inbounds [4 x i64], ptr %in, i64 0, i64 2
+ %3 = load i64, ptr %arrayidx.2, align 8
+ %arrayidx.3 = getelementptr inbounds [4 x i64], ptr %in, i64 0, i64 3
+ %4 = load i64, ptr %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
+}
diff --git a/llvm/test/tools/UpdateTestChecks/update_llc_test_checks/nvptx-basic.test b/llvm/test/tools/UpdateTestChecks/update_llc_test_checks/nvptx-basic.test
new file mode 100644
index 0000000000000..5bd29bd2679af
--- /dev/null
+++ b/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
diff --git a/llvm/utils/UpdateTestChecks/asm.py b/llvm/utils/UpdateTestChecks/asm.py
index bb594846ff3e8..01a25a6f52aa1 100644
--- a/llvm/utils/UpdateTestChecks/asm.py
+++ b/llvm/utils/UpdateTestChecks/asm.py
@@ -178,6 +178,31 @@ class string:
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\})?)? = .*)$',
@@ -388,6 +413,20 @@ def scrub_asm_csky(asm, args):
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),
@@ -426,6 +465,7 @@ def get_run_handler(triple):
'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 = ''
@@ -444,7 +484,7 @@ def get_run_handler(triple):
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,
diff --git a/llvm/utils/UpdateTestChecks/common.py b/llvm/utils/UpdateTestChecks/common.py
index 9745e60bc45ea..7ff3e3a183388 100644
--- a/llvm/utils/UpdateTestChecks/common.py
+++ b/llvm/utils/UpdateTestChecks/common.py
@@ -402,11 +402,12 @@ def do_scrub(body, scrubber, scrubber_args, extra):
# 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):
@@ -484,6 +485,15 @@ def process_run_line(self, function_re, scrubber, raw_tool_output, prefixes, is_
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
@@ -558,7 +568,7 @@ def process_run_line(self, function_re, scrubber, raw_tool_output, prefixes, is_
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):
@@ -835,11 +845,12 @@ def add_checks(output_lines, comment_marker, prefix_list, func_dict, func_name,
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.
@@ -912,13 +923,13 @@ def add_ir_checks(output_lines, comment_marker, prefix_list, func_dict,
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,
diff --git a/llvm/utils/UpdateTestChecks/isel.py b/llvm/utils/UpdateTestChecks/isel.py
index a77800a79bffa..41f6cf45008e5 100644
--- a/llvm/utils/UpdateTestChecks/isel.py
+++ b/llvm/utils/UpdateTestChecks/isel.py
@@ -50,7 +50,7 @@ def get_run_handler(triple):
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,
More information about the llvm-commits
mailing list