yaxunl created this revision.
yaxunl added reviewers: tra, rjmccall.
Herald added a subscriber: aprantl.
Herald added a reviewer: aaron.ballman.
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 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.
https://reviews.llvm.org/D79866
Files:
clang/lib/Sema/SemaDeclAttr.cpp
clang/test/CodeGenCUDA/kernel-dbg-info.cu
Index: clang/test/CodeGenCUDA/kernel-dbg-info.cu
===================================================================
--- /dev/null
+++ 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);
+}
Index: clang/lib/Sema/SemaDeclAttr.cpp
===================================================================
--- clang/lib/Sema/SemaDeclAttr.cpp
+++ clang/lib/Sema/SemaDeclAttr.cpp
@@ -4361,6 +4361,12 @@
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.CUDAIsDevice)
+ D->addAttr(NoDebugAttr::CreateImplicit(S.Context));
}
static void handleGNUInlineAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
Index: clang/test/CodeGenCUDA/kernel-dbg-info.cu
===================================================================
--- /dev/null
+++ 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);
+}
Index: clang/lib/Sema/SemaDeclAttr.cpp
===================================================================
--- clang/lib/Sema/SemaDeclAttr.cpp
+++ clang/lib/Sema/SemaDeclAttr.cpp
@@ -4361,6 +4361,12 @@
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.CUDAIsDevice)
+ D->addAttr(NoDebugAttr::CreateImplicit(S.Context));
}
static void handleGNUInlineAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
