yaxunl updated this revision to Diff 263831.
yaxunl retitled this revision from "[CUDA][HIP] Do not emit debug info for stub
function" to "[HIP] Do not emit debug info for stub function".
yaxunl added a comment.
limit change to HIP
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D79866/new/
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
@@ -4353,6 +4353,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.HIP && !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
@@ -4353,6 +4353,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.HIP && !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
