[PATCH] D41827: [DEBUG] Initial adaptation of NVPTX target for debug info emission.

Tue Apr 17 10:08:12 PDT 2018

ABataev added a comment.

In https://reviews.llvm.org/D41827#1069938, @echristo wrote:

> Couple of small comment then this is going to be ok.
>
> One question: "What work needs to happen to uncomment the debug info output?" In other words, what isn't done right now?
>
> -eric

The main problem is that the emitted debug info is correct syntactically, but not semantically. And I'm afraid that this may cause some problems with the cuda tools.

================
Comment at: lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp:39
+  const MCObjectFileInfo *FI = getStreamer().getContext().getObjectFileInfo();
+  // Emit closing brace for DWARF sections only.
+  if (CurSection && CurSection != FI->getTextSection() &&
----------------
echristo wrote:
> ABataev wrote:
> > echristo wrote:
> > > This should be commented as to "only dwarf sections need comments according to 'some reason'"
> > Ok
> For the record this is pretty fragile and should be fixed up in a followup patch.
Ok, will do

================
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
----------------
echristo wrote:
> ABataev wrote:
> > echristo wrote:
> > > 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?
> > Not for ptxas, ptxas supports this order of the `.loc` and `.file` directives.
> Can you document this in comments?
Sure

Repository:
  rL LLVM

https://reviews.llvm.org/D41827