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,31 @@ 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[^\(\n]+)' + + # function name separator (opening brace) + r'(?P\()' + + # 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,6 +413,20 @@ 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 @@ '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 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, Index: llvm/utils/UpdateTestChecks/common.py =================================================================== --- llvm/utils/UpdateTestChecks/common.py +++ llvm/utils/UpdateTestChecks/common.py @@ -402,11 +402,12 @@ # 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 @@ 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 @@ 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 @@ 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 @@ 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, Index: llvm/utils/UpdateTestChecks/isel.py =================================================================== --- llvm/utils/UpdateTestChecks/isel.py +++ llvm/utils/UpdateTestChecks/isel.py @@ -50,7 +50,7 @@ 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,