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

Wed Apr 11 08:14:56 PDT 2018

ABataev added inline comments.

================
Comment at: lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.cpp:42
+  // FIXME: remove comment once debug info is properly supported.
+  Data8bitsDirective = "// .b8 ";
+  Data16bitsDirective = nullptr; // not supported
----------------
echristo wrote:
> Do we still need these parts commented out?
Yes, cause the debug info is still not correct. I'm afraid that it may cause some problems in the ptxas

================
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:
> This should be commented as to "only dwarf sections need comments according to 'some reason'"
Ok

================
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:
> echristo wrote:
> > This needs more elaboration.
> Similar to my comment below, I think that the "which section am I in" part needs some work. I definitely don't think it should be a "!=" comparison, and perhaps should be checking for some other metadata/feature.
1. Do you suggest explicitly check for the particular debug sections? 
2. What metadata/feature are you talking about? Could you give a reference?

================
Comment at: lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp:43
+    OS << "//\t}\n";
+  if (Section != FI->getTextSection() && Section != FI->getDataSection() &&
+      Section != FI->getBSSSection()) {
----------------
echristo wrote:
> ABataev wrote:
> > echristo wrote:
> > > I don't understand what's going on here?
> > The body of the DWARF sections in PTX file must be enclosed into braces, like this:
> > ```
> > .section .debug_info
> > {
> > ...
> > }
> > ```
> > This code checks that current section is one of the DWARF sections and encloses the content of these sections into braces.
> But why looking for particular sections this way? 
Do you have some better solution in mind?

================
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:
> > > 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`  

================
Comment at: lib/Target/NVPTX/NVPTXAsmPrinter.h:350
+    bool Result = AsmPrinter::runOnMachineFunction(F);
+    OutStreamer->EmitRawText(StringRef("}\n"));
+    return Result;
----------------
echristo wrote:
> ABataev wrote:
> > tra wrote:
> > > Do I understand it correctly that printout of the `}` was moved here out of `EmitFunctionBodyEnd()` so that we can emit additional debug labels after the last basic block?
> > > 
> > > It would be good to add few comments describing why we emit (or not) braces in particular places.
> > Yes, correct.
> > Ok, will add
> Can you add a rationale here as to why we do it here rather than some other place? More detail in comment descriptions is good.
Ok, will add

Repository:
  rL LLVM

https://reviews.llvm.org/D41827