mirror of
https://github.com/compiler-explorer/compiler-explorer.git
synced 2026-09-11 10:12:09 -04:00
8040a92032
Fixes #8964. Directive filtering runs over the whole listing, including the region between `#APP` and `#NO_APP`. That region is the user's own inline asm, so filtering it changes what the block says. For the case in the issue: ```cpp int square(int num) { asm(".rept 5\nnop\n.endr"); return num * num; } ``` `.rept 5` and `.endr` go away and a single `nop` is left, which reads as one instruction where the source asks for five. The parser already tracks the region: `updateParsingState` calls `enterCustomAssembly` on `#APP` and `exitCustomAssembly` on `#NO_APP`, and `isInCustomAssembly()` was being used only to fix label indentation. `shouldSkipDirective` now consults it too. ### What still gets dropped Not everything in the block is the user's. GCC interleaves its own line markers, as @Alcaro pointed out on the issue: ``` #APP # 3 "/app/example.cpp" 1 .rept 5 nop .endr # 0 "" 2 .loc 1 4 5 view .LVU3 # /app/example.cpp:4: return num * num; .loc 1 4 18 is_stmt 0 view .LVU4 #NO_APP ``` So the rule is narrower than "skip filtering inside the block": inside it, a directive is dropped only when it is a source-location marker. That set is not a new list, it is the one `SourceLineHandler` already consumes to build the source mapping (`.loc`, `.d2line`, `.cv_loc`, `.dbg line`, `.stabn`), exposed as `isSourceLocationDirective`. Anything the handler does not claim is text the compiler was asked to assemble, and it stays. The `#` lines above are comments and are unaffected, they still go with the comment filter. ### The changed snapshots `nvcc-x86-host-example` gains three lines: `.section .nv_fatbin, "a"`, `.align 8` and `.text`, all inside its `#APP` block. The hundreds of `.quad` lines in that block were already surviving the directive filter, because `fatbinData:` sets `prevLabel` and `dataDefn` lines under a used label are kept, so this is three lines added to output that was already showing the blob, not the blob becoming visible. `NvccCompiler.removeNvccFatbinaryBlob` is what actually hides it, and it is unaffected. No other case in `filters-cases` moves. `bug-577_gcc` and `bug-577_clang` have `#APP` blocks with no directives in them. ### Tests Three in `test/asm-parser-tests.ts`: the directives survive, the `.loc` markers inside the same block still do not, and filtering resumes after `#NO_APP`. The first and third fail on `main`. I ran the full suite in `node:22-bookworm`, since `filter-tests.ts` returns early on Windows: 2631 passed.
619 lines
23 KiB
TypeScript
619 lines
23 KiB
TypeScript
// Copyright (c) 2024, Compiler Explorer Authors
|
|
// All rights reserved.
|
|
//
|
|
// Redistribution and use in source and binary forms, with or without
|
|
// modification, are permitted provided that the following conditions are met:
|
|
//
|
|
// * Redistributions of source code must retain the above copyright notice,
|
|
// this list of conditions and the following disclaimer.
|
|
// * Redistributions in binary form must reproduce the above copyright
|
|
// notice, this list of conditions and the following disclaimer in the
|
|
// documentation and/or other materials provided with the distribution.
|
|
//
|
|
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
|
|
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
|
|
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
|
|
// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
|
|
// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
|
|
// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
|
|
// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
|
|
// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
|
|
// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
|
|
// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
|
|
// POSSIBILITY OF SUCH DAMAGE.
|
|
|
|
import {describe, expect, it} from 'vitest';
|
|
|
|
import {AsmParser} from '../lib/parsers/asm-parser.js';
|
|
import {MlirAsmParser} from '../lib/parsers/asm-parser-mlir.js';
|
|
import {PolkaVMAsmParser} from '../lib/parsers/asm-parser-polkavm.js';
|
|
import {PTXAsmParser} from '../lib/parsers/asm-parser-ptx.js';
|
|
import {ResolcRiscVAsmParser} from '../lib/parsers/asm-parser-resolc-riscv.js';
|
|
import type {ParsedAsmResult} from '../types/asmresult/asmresult.interfaces.js';
|
|
import type {ParseFiltersAndOutputOptions} from '../types/features/filters.interfaces.js';
|
|
|
|
describe('AsmParser tests', () => {
|
|
const parser = new AsmParser();
|
|
it('should identify generic opcodes', () => {
|
|
expect(parser.hasOpcode(' mov r0, #1')).toBe(true);
|
|
expect(parser.hasOpcode(' ROL A')).toBe(true);
|
|
});
|
|
it('should not identify non-opcodes as opcodes', () => {
|
|
expect(parser.hasOpcode(' ;mov r0, #1')).toBe(false);
|
|
expect(parser.hasOpcode('')).toBe(false);
|
|
expect(parser.hasOpcode('# moose')).toBe(false);
|
|
});
|
|
it('should identify llvm opcodes', () => {
|
|
expect(parser.hasOpcode(' %i1 = phi i32 [ %i2, %.preheader ], [ 0, %bb ]')).toBe(true);
|
|
});
|
|
});
|
|
|
|
describe('AsmParser comment filtering', () => {
|
|
const parser = new AsmParser();
|
|
it('should keep label lines starting with @ when filtering comments', () => {
|
|
const input = '@cube@4:\n ret';
|
|
const result = parser.processAsm(input, {commentOnly: true});
|
|
const lines = result.asm.map(line => line.text);
|
|
expect(lines[0]).toBe('@cube@4:');
|
|
expect(lines[1]).toBe(' ret');
|
|
});
|
|
});
|
|
|
|
describe('AsmParser directive filtering', () => {
|
|
const parser = new AsmParser();
|
|
// GCC wraps inline asm in #APP/#NO_APP and interleaves its own .loc markers there.
|
|
const inlineAsm = `square(int):
|
|
pushq %rbp
|
|
#APP
|
|
# 3 "/app/example.cpp" 1
|
|
.rept 5
|
|
nop
|
|
.endr
|
|
# 0 "" 2
|
|
.loc 1 4 5 view .LVU3
|
|
#NO_APP
|
|
movl %edi, %eax
|
|
.loc 1 5 1
|
|
popq %rbp
|
|
ret`;
|
|
|
|
it('should keep inline asm directives when filtering directives', () => {
|
|
const result = parser.processAsm(inlineAsm, {directives: true});
|
|
const lines = result.asm.map(line => line.text.trim());
|
|
expect(lines).toContain('.rept 5');
|
|
expect(lines).toContain('.endr');
|
|
});
|
|
|
|
it('should still drop source location markers inside an inline asm block', () => {
|
|
const result = parser.processAsm(inlineAsm, {directives: true});
|
|
const lines = result.asm.map(line => line.text.trim());
|
|
expect(lines.some(line => line.startsWith('.loc'))).toBe(false);
|
|
});
|
|
|
|
it('should keep filtering directives once the inline asm block has ended', () => {
|
|
const result = parser.processAsm(
|
|
`#APP
|
|
.rept 5
|
|
#NO_APP
|
|
.p2align 4
|
|
ret`,
|
|
{directives: true},
|
|
);
|
|
const lines = result.asm.map(line => line.text.trim());
|
|
expect(lines).toContain('.rept 5');
|
|
expect(lines).not.toContain('.p2align 4');
|
|
});
|
|
});
|
|
|
|
describe('PTXAsmParser tests', () => {
|
|
const parser = new PTXAsmParser();
|
|
|
|
describe('Identifying opcodes', () => {
|
|
it('should identify regular PTX opcodes', () => {
|
|
expect(parser.hasOpcode(' mov.u32 %r25, %ctaid.x;')).toBe(true);
|
|
expect(parser.hasOpcode(' ld.global.v4.b32 { %r1, %r2, %r3, %r4 }, [ %rd1 + 0 ];')).toBe(true);
|
|
expect(parser.hasOpcode(' mul.wide.s32 %rd10, %r31, 4;')).toBe(true);
|
|
});
|
|
it('should identify PTX opcodes with predicate', () => {
|
|
expect(parser.hasOpcode(' @%p1 ld.global.v4.b32 { %r1, %r2, %r3, %r4 }, [ %rd1 + 0 ];')).toBe(true);
|
|
expect(parser.hasOpcode(' @!%p1 bra LBB6_2;')).toBe(true);
|
|
expect(parser.hasOpcode(' @%r789 bra LBB6_2;')).toBe(true);
|
|
});
|
|
it('should identify PTX opcodes wrapped in braces', () => {
|
|
expect(parser.hasOpcode('{fma.rn.f16x2 %r9,%r2,%r3,%r270;')).toBe(true);
|
|
expect(parser.hasOpcode(' {fma.rn.f16x2 %r9,%r2,%r3,%r270;')).toBe(true);
|
|
expect(parser.hasOpcode('{ add.s32 %r1, %r2, %r3;')).toBe(true);
|
|
});
|
|
});
|
|
|
|
describe('Nested brace indentation', () => {
|
|
it('should indent content inside callseq blocks', () => {
|
|
const input = `{ // callseq 0, 0
|
|
.param .b64 param0;
|
|
st.param.b64 [param0+0], %rd2;
|
|
}`;
|
|
const result = parser.processAsm(input, {});
|
|
const lines = result.asm.map(line => line.text);
|
|
|
|
expect(lines[0]).toBe('{ // callseq 0, 0');
|
|
expect(lines[1]).toBe('\t.param .b64 param0;');
|
|
expect(lines[2]).toBe('\tst.param.b64 [param0+0], %rd2;');
|
|
expect(lines[3]).toBe('}');
|
|
});
|
|
|
|
it('should properly indent function calls inside nested braces', () => {
|
|
const input = `{ // callseq 0, 0
|
|
call.uni (retval0),
|
|
vprintf,
|
|
(
|
|
param0,
|
|
param1
|
|
);
|
|
}`;
|
|
const result = parser.processAsm(input, {});
|
|
const lines = result.asm.map(line => line.text);
|
|
|
|
expect(lines[0]).toBe('{ // callseq 0, 0');
|
|
expect(lines[1]).toBe('\tcall.uni (retval0),');
|
|
expect(lines[2]).toBe('\tvprintf,');
|
|
expect(lines[3]).toBe('\t(');
|
|
expect(lines[4]).toBe('\tparam0,');
|
|
expect(lines[5]).toBe('\tparam1');
|
|
expect(lines[6]).toBe('\t);');
|
|
expect(lines[7]).toBe('}');
|
|
});
|
|
|
|
it('should have proper indentation logic', () => {
|
|
const input = `.visible .entry _Z6kernelv()
|
|
{
|
|
mov.u64 %rd1, $str;
|
|
{ // callseq 0, 0
|
|
call.uni (retval0),
|
|
vprintf,
|
|
(
|
|
param0,
|
|
param1
|
|
);
|
|
}
|
|
}`;
|
|
const result = parser.processAsm(input, {});
|
|
const lines = result.asm.map(line => line.text);
|
|
|
|
expect(lines[0]).toBe('.visible .entry _Z6kernelv()');
|
|
expect(lines[1]).toBe('{');
|
|
expect(lines[2]).toBe('\tmov.u64 %rd1, $str;');
|
|
expect(lines[3]).toBe('\t{ // callseq 0, 0');
|
|
expect(lines[4]).toBe('\t\tcall.uni (retval0),');
|
|
expect(lines[5]).toBe('\t\tvprintf,');
|
|
expect(lines[6]).toBe('\t\t(');
|
|
expect(lines[7]).toBe('\t\tparam0,');
|
|
expect(lines[8]).toBe('\t\tparam1');
|
|
expect(lines[9]).toBe('\t\t);');
|
|
expect(lines[10]).toBe('\t}');
|
|
expect(lines[11]).toBe('}');
|
|
});
|
|
});
|
|
|
|
describe('Label handling', () => {
|
|
it('should never indent labels', () => {
|
|
const input = `{ // callseq 0, 0
|
|
$L__BB0_3:
|
|
mov.u64 %rd1, $str;
|
|
}`;
|
|
const result = parser.processAsm(input, {});
|
|
const lines = result.asm.map(line => line.text);
|
|
|
|
expect(lines[0]).toBe('{ // callseq 0, 0');
|
|
expect(lines[1]).toBe('$L__BB0_3:');
|
|
expect(lines[2]).toBe('\tmov.u64 %rd1, $str;');
|
|
expect(lines[3]).toBe('}');
|
|
});
|
|
});
|
|
|
|
describe('Braced instructions', () => {
|
|
it('should not drop FMA instructions wrapped in braces on separate lines', () => {
|
|
const input = `{fma.rn.f16x2 %r9,%r2,%r3,%r270;
|
|
}
|
|
{fma.rn.f16x2 %r13,%r2,%r3,%r9;
|
|
}`;
|
|
const result = parser.processAsm(input, {});
|
|
const lines = result.asm.map(line => line.text);
|
|
|
|
expect(lines.length).toBe(4);
|
|
expect(lines[0]).toBe('{fma.rn.f16x2 %r9,%r2,%r3,%r270;');
|
|
expect(lines[1]).toBe('}');
|
|
expect(lines[2]).toBe('{fma.rn.f16x2 %r13,%r2,%r3,%r9;');
|
|
expect(lines[3]).toBe('}');
|
|
});
|
|
|
|
it('should preserve braced instructions inside functions', () => {
|
|
const input = `.visible .entry kernel()
|
|
{
|
|
mov.u32 %r1, 0;
|
|
{fma.rn.f16x2 %r9,%r2,%r3,%r270;
|
|
}
|
|
{fma.rn.f16x2 %r13,%r2,%r3,%r9;
|
|
}
|
|
ret;
|
|
}`;
|
|
const result = parser.processAsm(input, {});
|
|
const lines = result.asm.map(line => line.text);
|
|
|
|
expect(lines[0]).toBe('.visible .entry kernel()');
|
|
expect(lines[1]).toBe('{');
|
|
expect(lines[2]).toBe('\tmov.u32 %r1, 0;');
|
|
expect(lines[3]).toBe('\t{fma.rn.f16x2 %r9,%r2,%r3,%r270;');
|
|
expect(lines[4]).toBe('\t}');
|
|
expect(lines[5]).toBe('\t{fma.rn.f16x2 %r13,%r2,%r3,%r9;');
|
|
expect(lines[6]).toBe('\t}');
|
|
expect(lines[7]).toBe('\tret;');
|
|
expect(lines[8]).toBe('}');
|
|
});
|
|
});
|
|
|
|
describe('Trim filter', () => {
|
|
it('should convert tabs to spaces when trim filter is enabled', () => {
|
|
const input = `\t\t\tcall.uni (retval0),
|
|
\t\t\tvprintf,
|
|
\t\t\t(
|
|
\t\t\tparam0,
|
|
\t\t\tparam1
|
|
\t\t\t);`;
|
|
const result = parser.processAsm(input, {trim: true});
|
|
const lines = result.asm.map(line => line.text);
|
|
|
|
expect(lines[0]).toBe(' call.uni (retval0),');
|
|
expect(lines[1]).toBe(' vprintf,');
|
|
expect(lines[2]).toBe(' (');
|
|
expect(lines[3]).toBe(' param0,');
|
|
expect(lines[4]).toBe(' param1');
|
|
expect(lines[5]).toBe(' );');
|
|
});
|
|
|
|
it('should preserve structure but compress whitespace with trim filter', () => {
|
|
const input = `{ // callseq 0, 0
|
|
\t\t\tcall.uni (retval0),
|
|
\t\t\tvprintf,
|
|
}`;
|
|
const result = parser.processAsm(input, {trim: true});
|
|
const lines = result.asm.map(line => line.text);
|
|
|
|
expect(lines[0]).toBe('{ // callseq 0, 0');
|
|
expect(lines[1]).toBe(' call.uni (retval0),');
|
|
expect(lines[2]).toBe(' vprintf,');
|
|
expect(lines[3]).toBe('}');
|
|
});
|
|
|
|
it('should preserve nested indentation structure with trim filter', () => {
|
|
const input = `.visible .entry _Z6kernelv()
|
|
{
|
|
mov.u64 %rd1, $str;
|
|
{ // callseq 0, 0
|
|
call.uni (retval0),
|
|
vprintf,
|
|
}
|
|
}`;
|
|
const result = parser.processAsm(input, {trim: true});
|
|
const lines = result.asm.map(line => line.text);
|
|
|
|
expect(lines[0]).toBe('.visible .entry _Z6kernelv()');
|
|
expect(lines[1]).toBe('{');
|
|
expect(lines[2]).toBe(' mov.u64 %rd1, $str;');
|
|
expect(lines[3]).toBe(' { // callseq 0, 0');
|
|
expect(lines[4]).toBe(' call.uni (retval0),');
|
|
expect(lines[5]).toBe(' vprintf,');
|
|
expect(lines[6]).toBe(' }');
|
|
expect(lines[7]).toBe('}');
|
|
});
|
|
});
|
|
|
|
describe('Directive filter', () => {
|
|
it('should remove directives when directive filter is enabled', () => {
|
|
const input = `
|
|
.file 1 "/tmp/compiler-explorer-compilerk2brih/example.py"
|
|
.section .debug_abbrev
|
|
{
|
|
.b8 1
|
|
}
|
|
.section .debug_info
|
|
{
|
|
.b32 70
|
|
.b32 .debug_abbrev
|
|
}
|
|
.section .debug_macinfo { }
|
|
`;
|
|
const result = parser.processAsm(input, {directives: true});
|
|
const lines = result.asm.map(line => line.text);
|
|
expect(lines.length).toBe(1);
|
|
expect(lines[0]).toBe('');
|
|
});
|
|
});
|
|
});
|
|
|
|
describe('MlirAsmParser tests', () => {
|
|
const parser = new MlirAsmParser();
|
|
|
|
describe('Location handling', () => {
|
|
it('should process MLIR with location references', () => {
|
|
const input = `
|
|
#loc = loc("<source>":7:0)
|
|
module {
|
|
tt.func public @add_kernel(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("<source>":7:0), %arg1: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("<source>":7:0), %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("<source>":7:0), %arg3: i32 {tt.divisibility = 16 : i32} loc("<source>":7:0)) attributes {noinline = false} {
|
|
%c1024_i32 = arith.constant 1024 : i32 loc(#loc1)
|
|
%0 = tt.get_program_id x : i32 loc(#loc2)
|
|
} loc(#loc)
|
|
} loc(#loc)
|
|
#loc1 = loc(unknown)
|
|
#loc2 = loc("<source>":14:24)`;
|
|
|
|
const result = parser.processAsm(input, {});
|
|
|
|
// Verify location definitions are removed
|
|
expect(result.asm.find(line => line.text.includes('#loc = loc'))).toBeUndefined();
|
|
expect(result.asm.find(line => line.text.includes('#loc1 = loc'))).toBeUndefined();
|
|
expect(result.asm.find(line => line.text.includes('#loc2 = loc'))).toBeUndefined();
|
|
|
|
// Verify location references are removed from displayed text
|
|
expect(result.asm.find(line => line.text.includes('loc(#loc)'))).toBeUndefined();
|
|
expect(result.asm.find(line => line.text.includes('loc(#loc1)'))).toBeUndefined();
|
|
expect(result.asm.find(line => line.text.includes('loc(#loc2)'))).toBeUndefined();
|
|
|
|
// Verify inline locations are removed
|
|
expect(result.asm.find(line => line.text.includes('loc("<source>"'))).toBeUndefined();
|
|
|
|
// Verify source information is correctly associated
|
|
const programIdLine = result.asm.find(line => line.text.includes('tt.get_program_id'));
|
|
expect(programIdLine).toBeDefined();
|
|
expect(programIdLine?.source).toBeDefined();
|
|
expect(programIdLine?.source?.file).toBe('<source>');
|
|
expect(programIdLine?.source?.line).toBe(14);
|
|
expect(programIdLine?.source?.column).toBe(24);
|
|
|
|
// Verify unknown locations are not associated with source information
|
|
const constantLine = result.asm.find(line => line.text.includes('arith.constant'));
|
|
expect(constantLine).toBeDefined();
|
|
expect(constantLine?.source).toBeNull();
|
|
|
|
// Verify the structure is preserved
|
|
const moduleStartLine = result.asm.find(line => line.text.includes('module {'));
|
|
expect(moduleStartLine).toBeDefined();
|
|
|
|
const funcLine = result.asm.find(line => line.text.includes('tt.func public @add_kernel('));
|
|
expect(funcLine).toBeDefined();
|
|
expect(funcLine?.text?.includes('%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, ')).toBe(true);
|
|
expect(funcLine?.text?.includes('%arg1: !tt.ptr<f32> {tt.divisibility = 16 : i32}, ')).toBe(true);
|
|
expect(funcLine?.text?.includes('%arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32}, ')).toBe(true);
|
|
expect(funcLine?.text?.includes('%arg3: i32 {tt.divisibility = 16 : i32})')).toBe(true);
|
|
|
|
const constLine = result.asm.find(line => line.text.includes('arith.constant 1024'));
|
|
expect(constLine).toBeDefined();
|
|
});
|
|
|
|
it('should resolve nested named MLIR locations', () => {
|
|
const input = `#loc = loc("<source>":16:0)
|
|
#loc1 = loc("<source>":17:24)
|
|
#loc2 = loc("<source>":18:28)
|
|
#loc19 = loc("pid"(#loc1))
|
|
#loc20 = loc("block_start"(#loc2))
|
|
#loc21 = loc(unknown)
|
|
#loc22 = loc("unknown_value"(#loc21))
|
|
module {
|
|
tt.func public @kernel(%x_ptr: !tt.ptr<f32> loc("x_ptr"(#loc))) attributes {noinline = false} {
|
|
%pid = tt.get_program_id x : i32 loc(#loc19)
|
|
%block_start = arith.muli %pid, %c1024_i32 : i32 loc(#loc20)
|
|
%unknown = arith.constant 0 : i32 loc(#loc22)
|
|
} loc(#loc)
|
|
}`;
|
|
|
|
const result = parser.processAsm(input, {});
|
|
|
|
expect(result.asm[0].text).toBe('module {');
|
|
expect(result.asm.find(line => line.text.includes('#loc'))).toBeUndefined();
|
|
expect(result.asm.find(line => line.text.includes('loc('))).toBeUndefined();
|
|
|
|
const functionLine = result.asm.find(line => line.text.includes('tt.func public @kernel'));
|
|
expect(functionLine?.source?.line).toBe(16);
|
|
expect(functionLine?.source?.column).toBe(0);
|
|
|
|
const pidLine = result.asm.find(line => line.text.includes('%pid = tt.get_program_id'));
|
|
expect(pidLine?.source?.line).toBe(17);
|
|
expect(pidLine?.source?.column).toBe(24);
|
|
|
|
const blockStartLine = result.asm.find(line => line.text.includes('%block_start = arith.muli'));
|
|
expect(blockStartLine?.source?.line).toBe(18);
|
|
expect(blockStartLine?.source?.column).toBe(28);
|
|
|
|
const unknownLine = result.asm.find(line => line.text.includes('%unknown = arith.constant'));
|
|
expect(unknownLine?.source).toBeNull();
|
|
});
|
|
});
|
|
});
|
|
|
|
describe('ResolcRiscVAsmParser tests', () => {
|
|
const parser = new ResolcRiscVAsmParser();
|
|
|
|
function expectParsedAsmResult(result: ParsedAsmResult, expected: ParsedAsmResult): void {
|
|
expect(result.labelDefinitions).toEqual(expected.labelDefinitions);
|
|
expect(result.asm.length).toEqual(expected.asm.length);
|
|
for (let i = 0; i < result.asm.length; i++) {
|
|
expect(result.asm[i]).toMatchObject(expected.asm[i]);
|
|
}
|
|
}
|
|
|
|
it.skipIf(process.platform === 'win32')('should identify RISC-V instruction info and source line numbers', () => {
|
|
const filters: Partial<ParseFiltersAndOutputOptions> = {binaryObject: true};
|
|
const riscv = `
|
|
000000000000027a <__entry>:
|
|
; __entry():
|
|
; path/to/example.sol.Square.yul:1
|
|
27a: 41 11 addi sp, sp, -0x10
|
|
|
|
000000000000028c <.Lpcrel_hi4>:
|
|
; .Lpcrel_hi4():
|
|
28c: 97 05 00 00 auipc a1, 0x0
|
|
; path/to/example.sol.Square.yul:1
|
|
2a8: 1d 71 addi sp, sp, -0x60
|
|
2aa: 86 ec sd ra, 0x58(sp)
|
|
; path/to/example.sol.Square.yul:7
|
|
2ca: e7 80 00 00 jalr ra <.Lpcrel_hi4+0x3a>`;
|
|
|
|
const expected: ParsedAsmResult = {
|
|
asm: [
|
|
{
|
|
text: '__entry:',
|
|
source: null,
|
|
},
|
|
{
|
|
text: ' addi sp, sp, -0x10',
|
|
address: 0x27a,
|
|
opcodes: ['41', '11'],
|
|
source: {
|
|
line: 1,
|
|
file: null,
|
|
},
|
|
},
|
|
{
|
|
text: '.Lpcrel_hi4:',
|
|
source: null,
|
|
},
|
|
{
|
|
text: ' auipc a1, 0x0',
|
|
address: 0x28c,
|
|
opcodes: ['97', '05', '00', '00'],
|
|
source: {
|
|
line: 1,
|
|
file: null,
|
|
},
|
|
},
|
|
{
|
|
text: ' addi sp, sp, -0x60',
|
|
address: 0x2a8,
|
|
opcodes: ['1d', '71'],
|
|
source: {
|
|
line: 1,
|
|
file: null,
|
|
},
|
|
},
|
|
{
|
|
text: ' sd ra, 0x58(sp)',
|
|
address: 0x2aa,
|
|
opcodes: ['86', 'ec'],
|
|
source: {
|
|
line: 1,
|
|
file: null,
|
|
},
|
|
},
|
|
{
|
|
text: ' jalr ra <.Lpcrel_hi4+0x3a>',
|
|
address: 0x2ca,
|
|
opcodes: ['e7', '80', '00', '00'],
|
|
source: {
|
|
line: 7,
|
|
file: null,
|
|
},
|
|
},
|
|
],
|
|
labelDefinitions: {
|
|
__entry: 1,
|
|
['.Lpcrel_hi4']: 3,
|
|
},
|
|
};
|
|
|
|
const result = parser.processAsm(riscv, filters);
|
|
expectParsedAsmResult(result, expected);
|
|
});
|
|
});
|
|
|
|
describe('PolkaVMAsmParser tests', () => {
|
|
const parser = new PolkaVMAsmParser();
|
|
|
|
function expectParsedAsmResult(result: ParsedAsmResult, expected: ParsedAsmResult): void {
|
|
expect(result.labelDefinitions).toEqual(expected.labelDefinitions);
|
|
expect(result.asm.length).toEqual(expected.asm.length);
|
|
for (let i = 0; i < result.asm.length; i++) {
|
|
expect(result.asm[i]).toMatchObject(expected.asm[i]);
|
|
}
|
|
}
|
|
|
|
// Note: We currently have no source mappings from PVM.
|
|
it('should identify PVM instruction info', () => {
|
|
const filters: Partial<ParseFiltersAndOutputOptions> = {
|
|
binaryObject: false,
|
|
commentOnly: false,
|
|
};
|
|
const pvm = `
|
|
// Code size = 1078 bytes
|
|
|
|
<__entry>:
|
|
: @0 (gas: 6)
|
|
0: sp = sp + 0xfffffffffffffff0
|
|
3: u64 [sp + 0x8] = ra
|
|
6: u64 [sp] = s0
|
|
8: s0 = a0 & 0x1
|
|
11: ecalli 2 // 'call_data_size'
|
|
13: fallthrough
|
|
: @1 (gas: 2)
|
|
14: u32 [0x20000] = a0`;
|
|
|
|
const expected: ParsedAsmResult = {
|
|
asm: [
|
|
{
|
|
text: '// Code size = 1078 bytes',
|
|
source: null,
|
|
},
|
|
{
|
|
text: '__entry:',
|
|
source: null,
|
|
},
|
|
{
|
|
text: ' @0 (gas: 6)',
|
|
source: null,
|
|
},
|
|
{
|
|
text: ' sp = sp + 0xfffffffffffffff0',
|
|
address: 0,
|
|
source: null,
|
|
},
|
|
{
|
|
text: ' u64 [sp + 0x8] = ra',
|
|
address: 3,
|
|
source: null,
|
|
},
|
|
{
|
|
text: ' u64 [sp] = s0',
|
|
address: 6,
|
|
source: null,
|
|
},
|
|
{
|
|
text: ' s0 = a0 & 0x1',
|
|
address: 8,
|
|
source: null,
|
|
},
|
|
{
|
|
text: " ecalli 2 // 'call_data_size'",
|
|
address: 11,
|
|
source: null,
|
|
},
|
|
{
|
|
text: ' fallthrough',
|
|
address: 13,
|
|
source: null,
|
|
},
|
|
{
|
|
text: ' @1 (gas: 2)',
|
|
source: null,
|
|
},
|
|
{
|
|
text: ' u32 [0x20000] = a0',
|
|
address: 14,
|
|
source: null,
|
|
},
|
|
],
|
|
labelDefinitions: {__entry: 2},
|
|
};
|
|
|
|
const result = parser.process(pvm, filters);
|
|
expectParsedAsmResult(result, expected);
|
|
});
|
|
});
|