wangtao
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

[HIP] Do not emit debug info for stub function 

The stub function is generated by compiler and its instructions have nothing
to do with the kernel source code.

Currently clang generates debug info for the stub function, which causes
confusion for the HIP debugger. For example, when users set break point
on a line of a kernel, the debugger should break on that line when the kernel is
executed and reaches that line, but instead the debugger breaks in the stub function.

This patch disables debug info for stub function for HIP.

Differential Revision: https://reviews.llvm.org/D79866
This commit is contained in:
Yaxun (Sam) Liu 2020-05-13 09:02:30 -04:00
parent 591b029f40
commit 1b7bf1bd75
2 changed files with 39 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
clang/lib/Sema/SemaDeclAttr.cpp
View File

@ -4361,6 +4361,12 @@ static void handleGlobalAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
S.Diag(FD->getBeginLoc(), diag::warn_kern_is_inline) << FD;
D->addAttr(::new (S.Context) CUDAGlobalAttr(S.Context, AL));
// In host compilation the kernel is emitted as a stub function, which is
// a helper function for launching the kernel. The instructions in the helper
// function has nothing to do with the source code of the kernel. Do not emit
// debug info for the stub function to avoid confusing the debugger.
if (S.LangOpts.HIP && !S.LangOpts.CUDAIsDevice)
D->addAttr(NoDebugAttr::CreateImplicit(S.Context));
}
static void handleGNUInlineAttr(Sema &S, Decl *D, const ParsedAttr &AL) {

33
clang/test/CodeGenCUDA/kernel-dbg-info.cu Normal file
View File

@ -0,0 +1,33 @@
// RUN: echo "GPU binary would be here" > %t
// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -O0 \
// RUN: -fcuda-include-gpubinary %t -debug-info-kind=limited \
// RUN: -o - -x hip | FileCheck %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm %s -O0 \
// RUN: -fcuda-include-gpubinary %t -debug-info-kind=limited \
// RUN: -o - -x hip -fcuda-is-device | FileCheck -check-prefix=DEV %s
#include "Inputs/cuda.h"
extern "C" __global__ void ckernel(int *a) {
*a = 1;
}
// Device side kernel names
// CHECK: @[[CKERN:[0-9]*]] = {{.*}} c"ckernel\00"
// DEV: define {{.*}}@ckernel{{.*}}!dbg
// DEV: store {{.*}}!dbg
// DEV: ret {{.*}}!dbg
// CHECK-NOT: define {{.*}}@__device_stub__ckernel{{.*}}!dbg
// CHECK: define {{.*}}@[[CSTUB:__device_stub__ckernel]]
// CHECK-NOT: call {{.*}}@hipLaunchByPtr{{.*}}!dbg
// CHECK: call {{.*}}@hipLaunchByPtr{{.*}}@[[CSTUB]]
// CHECK-NOT: ret {{.*}}!dbg
// CHECK-LABEL: define {{.*}}@_Z8hostfuncPi{{.*}}!dbg
// CHECK: call void @[[CSTUB]]{{.*}}!dbg
void hostfunc(int *a) {
ckernel<<<1, 1>>>(a);
}