[PATCH] D41827: [DEBUG] Initial adaptation of NVPTX target for debug info emission.
Eric Christopher via Phabricator via llvm-commits
llvm-commits at lists.llvm.org
Wed Apr 11 10:23:00 PDT 2018
echristo added inline comments.
================
Comment at: lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h:33
+ /// same level as kernel and device function declarations.
+ /// LLVM emits .file directives immediately the location debug info is
+ /// emitted, i.e. they may be emitted inside functions. We gather all these
----------------
ABataev wrote:
> echristo wrote:
> > ABataev wrote:
> > > echristo wrote:
> > > > So, this functionality is used for changing files within the line table. I think what's actually going on here is that nvptx, for some reason, only supports a single file directive.
> > > >
> > > > Basically you don't support the idea of:
> > > >
> > > > int foo (void) {
> > > > // some code
> > > > #include "somefile.def"
> > > > // some more code
> > > > }
> > > >
> > > > having a file switch in the middle of a function. That's arguably a bug, but I think you don't want to collect them and emit them outside the functions because then it would be more actively wrong.
> > > No, it means that you may have the reference to the `somefile.def`, but the table of files should be emitted outside of the function, in the outer context
> > I'm not sure what you mean. Can you show some examples and elaborate? I.e. you may have a file table entry, but if you're not emitting them before the function you might have a bad time?
> >
> > What do you do with this testcase?
> >
> > echristo at athyra ~> cat foo.c
> > void foo() {
> > #include "foo.def"
> > return;
> > }
> > echristo at athyra ~> cat foo.def
> > #include <stdio.h>
> > printf("hello world\n");
> I removed `include <stdio.h>` from your second file, it causes a lot of troubles. Here is what I got:
> ```
> In file included from /NVPTX_DEBUG_INFO/foo.c:3:
> ./foo.def:2:1: warning: implicitly declaring library function 'printf' with type 'int (const char *, ...)' [-Wimplicit-function-declaration]
> printf("hello world\n");
> ^
> ./foo.def:2:1: note: include the header <stdio.h> or explicitly provide a declaration for 'printf'
> //
> // Generated by LLVM NVPTX Back-End
> //
>
> .version 3.2
> .target sm_20//, debug
> .address_size 64
>
> // .globl foo
> .extern .func (.param .b32 func_retval0) vprintf
> (
> .param .b64 vprintf_param_0,
> .param .b64 vprintf_param_1
> )
> ;
> .global .align 1 .b8 _$_str[13] = {104, 101, 108, 108, 111, 32, 119, 111, 114, 108, 100, 10, 0};
>
> .visible .func foo()
> {
> .reg .b32 %r<2>;
> .reg .b64 %rd<4>;
> Lfunc_begin0:
> .loc 1 1 0
>
> .loc 2 2 1
> mov.u64 %rd1, _$_str;
> cvta.global.u64 %rd2, %rd1;
> mov.u64 %rd3, 0;
> { // callseq 0, 0
> .reg .b32 temp_param_reg;
> .param .b64 param0;
> st.param.b64 [param0+0], %rd2;
> .param .b64 param1;
> st.param.b64 [param1+0], %rd3;
> .param .b32 retval0;
> call.uni (retval0),
> vprintf,
> (
> param0,
> param1
> );
> ld.param.b32 %r1, [retval0+0];
> } // callseq 0
> .loc 1 4 1
> ret;
> Ltmp0:
> Lfunc_end0:
>
> }
> .file 1 "/NVPTX_DEBUG_INFO/foo.c"
> .file 2 "/NVPTX_DEBUG_INFO/./foo.def"
> // .section .debug_abbrev
> // {
> // .b8 1
> // .b8 17
> // .b8 1
> // .b8 37
> // .b8 8
> // .b8 19
> // .b8 5
> // .b8 3
> // .b8 8
> // .b8 16
> // .b8 6
> // .b8 27
> // .b8 8
> // .b8 17
> // .b8 1
> // .b8 18
> // .b8 1
> // .b8 0
> // .b8 0
> // .b8 2
> // .b8 46
> // .b8 0
> // .b8 17
> // .b8 1
> // .b8 18
> // .b8 1
> // .b8 3
> // .b8 8
> // .b8 58
> // .b8 11
> // .b8 59
> // .b8 11
> // .b8 63
> // .b8 12
> // .b8 0
> // .b8 0
> // .b8 0
> // }
> // .section .debug_info
> // {
> // .b32 166
> // .b8 2
> // .b8 0
> // .b32 .debug_abbrev
> // .b8 8
> // .b8 1
> // .b8 99
> // .b8 108
> // .b8 97
> // .b8 110
> // .b8 103
> // .b8 32
> // .b8 118
> // .b8 101
> // .b8 114
> // .b8 115
> // .b8 105
> // .b8 111
> // .b8 110
> // .b8 32
> // .b8 55
> // .b8 46
> // .b8 48
> // .b8 46
> // .b8 48
> // .b8 32
> // .b8 40
> // .b8 116
> // .b8 114
> // .b8 117
> // .b8 110
> // .b8 107
> // .b8 32
> // .b8 51
> // .b8 50
> // .b8 57
> // .b8 56
> // .b8 48
> // .b8 52
> // .b8 41
> // .b8 32
> // .b8 40
> // .b8 108
> // .b8 108
> // .b8 118
> // .b8 109
> // .b8 47
> // .b8 116
> // .b8 114
> // .b8 117
> // .b8 110
> // .b8 107
> // .b8 32
> // .b8 51
> // .b8 50
> // .b8 57
> // .b8 56
> // .b8 48
> // .b8 53
> // .b8 41
> // .b8 0
> // .b8 12
> // .b8 0
> // .b8 102
> // .b8 111
> // .b8 111
> // .b8 46
> // .b8 99
> // .b8 0
> // .b32 .debug_line
> // .b8 47
> // .b8 103
> // .b8 115
> // .b8 97
> // .b8 47
> // .b8 121
> // .b8 107
> // .b8 116
> // .b8 103
> // .b8 115
> // .b8 97
> // .b8 47
> // .b8 104
> // .b8 111
> // .b8 109
> // .b8 101
> // .b8 47
> // .b8 97
> // .b8 47
> // .b8 108
> // .b8 47
> // .b8 97
> // .b8 108
> // .b8 101
> // .b8 120
> // .b8 101
> // .b8 121
> // .b8 47
> // .b8 116
> // .b8 101
> // .b8 115
> // .b8 116
> // .b8 47
> // .b8 78
> // .b8 86
> // .b8 80
> // .b8 84
> // .b8 88
> // .b8 95
> // .b8 68
> // .b8 69
> // .b8 66
> // .b8 85
> // .b8 71
> // .b8 95
> // .b8 73
> // .b8 78
> // .b8 70
> // .b8 79
> // .b8 0
> // .b64 Lfunc_begin0
> // .b64 Lfunc_end0
> // .b8 2
> // .b64 Lfunc_begin0
> // .b64 Lfunc_end0
> // .b8 102
> // .b8 111
> // .b8 111
> // .b8 0
> // .b8 1
> // .b8 1
> // .b8 1
> // .b8 0
> // }
> // .section .debug_macinfo
> // {
> // .b8 0
>
> // }
> 1 warning generated.
>
> ```
> As you can see `.loc` directives are emitted with references to files 1 and 2 and after the `.text` section we have 2 `.file` directives: one for file `foo.c` and another one for file `foo.def`
Typically you don't have file directives listed after the function and still have it able to assemble. Doing that in the .s file created with my testcase gives:
echristo at athyra ~> clang -c foo.s
foo.s:8:7: error: unassigned file number in '.loc' directive
.loc 1 1 0 # foo.c:1:0
^
foo.s:18:7: error: unassigned file number in '.loc' directive
.loc 2 2 1 prologue_end # ./foo.def:2:1
^
foo.s:22:7: error: unassigned file number in '.loc' directive
.loc 1 3 3 # foo.c:3:3
which seems problematic?
Repository:
rL LLVM
https://reviews.llvm.org/D41827
More information about the llvm-commits
mailing list