Skip to content

Commit

Permalink
[as] Add testcase to document that it reorders '.file' directives [#20]
Browse files Browse the repository at this point in the history
This relates to commit bd12880 "Deal with
.file and .loc directives".
  • Loading branch information
tschwinge committed Jan 7, 2021
1 parent b4e1cbf commit c34d2b5
Show file tree
Hide file tree
Showing 5 changed files with 123 additions and 0 deletions.
5 changes: 5 additions & 0 deletions test/as/debug/directive-file-1-sum.c
Original file line number Diff line number Diff line change
@@ -0,0 +1,5 @@
inline __attribute__((always_inline))
int sum(int a, int b)
{
return a + b;
}
6 changes: 6 additions & 0 deletions test/as/debug/directive-file-1.c
Original file line number Diff line number Diff line change
@@ -0,0 +1,6 @@
#include "directive-file-1-sum.c"

int f(int a, int b)
{
return sum(a, b);
}
49 changes: 49 additions & 0 deletions test/as/debug/directive-file-1.o.golden
Original file line number Diff line number Diff line change
@@ -0,0 +1,49 @@
// $ [nvptx-none-gcc] directive-file-1.c -o directive-file-1.s -S -g
// $ [nvptx-none-as] directive-file-1.s -o directive-file-1.o.golden
// BEGIN PREAMBLE
.version 3.1
.target sm_30
.address_size 64
// END PREAMBLE
// BEGIN GLOBAL FUNCTION DECL: f
.visible .func (.param .u32 %value_out) f (.param .u32 %in_ar0, .param .u32 %in_ar1);
.file 1 "directive-file-1-sum.c"
.file 2 "directive-file-1.c"
// BEGIN GLOBAL FUNCTION DEF: f
.visible .func (.param .u32 %value_out) f (.param .u32 %in_ar0, .param .u32 %in_ar1)
{
.reg .u32 %value;
.reg .u32 %ar0;
ld.param.u32 %ar0,[%in_ar0];
.reg .u32 %ar1;
ld.param.u32 %ar1,[%in_ar1];
.local .align 16 .b8 %frame_ar[24];
.reg .u64 %frame;
cvta.local.u64 %frame,%frame_ar;
.reg .u32 %r22;
.reg .u32 %r23;
.reg .u32 %r24;
.reg .u32 %r25;
.reg .u32 %r26;
.reg .u32 %r27;
.reg .u32 %r28;
.reg .u32 %r29;
mov.u32 %r24,%ar0;
st.u32 [%frame+16],%r24;
mov.u32 %r25,%ar1;
st.u32 [%frame+20],%r25;
ld.u32 %r26,[%frame+16];
st.u32 [%frame],%r26;
ld.u32 %r27,[%frame+20];
st.u32 [%frame+4],%r27;
.loc 1 4 12
ld.u32 %r28,[%frame];
ld.u32 %r29,[%frame+4];
add.u32 %r22,%r28,%r29;
.loc 2 5 10
mov.u32 %r23,%r22;
.loc 2 6 1
mov.u32 %value,%r23;
st.param.u32 [%value_out],%value;
ret;
}
54 changes: 54 additions & 0 deletions test/as/debug/directive-file-1.s
Original file line number Diff line number Diff line change
@@ -0,0 +1,54 @@
// $ [nvptx-none-gcc] directive-file-1.c -o directive-file-1.s -S -g
// $ [nvptx-none-as] directive-file-1.s -o directive-file-1.o.golden


// BEGIN PREAMBLE
.version 3.1
.target sm_30
.address_size 64
// END PREAMBLE


// BEGIN GLOBAL FUNCTION DECL: f
.visible .func (.param.u32 %value_out) f (.param.u32 %in_ar0, .param.u32 %in_ar1);

// BEGIN GLOBAL FUNCTION DEF: f
.visible .func (.param.u32 %value_out) f (.param.u32 %in_ar0, .param.u32 %in_ar1)
{
.reg.u32 %value;
.reg.u32 %ar0;
ld.param.u32 %ar0, [%in_ar0];
.reg.u32 %ar1;
ld.param.u32 %ar1, [%in_ar1];
.local .align 16 .b8 %frame_ar[24];
.reg.u64 %frame;
cvta.local.u64 %frame, %frame_ar;
.reg.u32 %r22;
.reg.u32 %r23;
.reg.u32 %r24;
.reg.u32 %r25;
.reg.u32 %r26;
.reg.u32 %r27;
.reg.u32 %r28;
.reg.u32 %r29;
mov.u32 %r24, %ar0;
st.u32 [%frame+16], %r24;
mov.u32 %r25, %ar1;
st.u32 [%frame+20], %r25;
ld.u32 %r26, [%frame+16];
st.u32 [%frame], %r26;
ld.u32 %r27, [%frame+20];
st.u32 [%frame+4], %r27;
.file 1 "directive-file-1-sum.c"
.loc 1 4 12
ld.u32 %r28, [%frame];
ld.u32 %r29, [%frame+4];
add.u32 %r22, %r28, %r29;
.file 2 "directive-file-1.c"
.loc 2 5 10
mov.u32 %r23, %r22;
.loc 2 6 1
mov.u32 %value, %r23;
st.param.u32 [%value_out], %value;
ret;
}
9 changes: 9 additions & 0 deletions test/as/debug/directive-file-1.test
Original file line number Diff line number Diff line change
@@ -0,0 +1,9 @@
<https://github.com/MentorEmbedded/nvptx-tools/issues/20> 'nvptx-as reorders ".file" directives'

In the file produced by GCC, expect 'nvptx-none-as' to reorder '.file' directives:
RUN: %target_as_cmd --no-verify %S/directive-file-1.s -o %t
RUN: cmp %S/directive-file-1.o.golden %t

... but expect further 'nvptx-none-as' runs to be idempotent:
RUN: %target_as_cmd --no-verify %S/directive-file-1.o.golden -o %t
RUN: cmp %S/directive-file-1.o.golden %t

0 comments on commit c34d2b5

Please sign in to comment.