[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