Author: Yaxun (Sam) Liu Date: 2020-05-13T17:55:40-04:00 New Revision: 1b7bf1bd75dcc4a362a83e4abc1200aaa75a38be
URL: https://github.com/llvm/llvm-project/commit/1b7bf1bd75dcc4a362a83e4abc1200aaa75a38be DIFF: https://github.com/llvm/llvm-project/commit/1b7bf1bd75dcc4a362a83e4abc1200aaa75a38be.diff LOG: [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 Added: clang/test/CodeGenCUDA/kernel-dbg-info.cu Modified: clang/lib/Sema/SemaDeclAttr.cpp Removed: ################################################################################ diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 883ab2ad0505..0e062ba74a08 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -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) { diff --git a/clang/test/CodeGenCUDA/kernel-dbg-info.cu b/clang/test/CodeGenCUDA/kernel-dbg-info.cu new file mode 100644 index 000000000000..a1a70d2cbaf2 --- /dev/null +++ b/clang/test/CodeGenCUDA/kernel-dbg-info.cu @@ -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); +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
