<div dir="ltr"><br><div class="gmail_extra"><br><div class="gmail_quote">On Mon, Dec 7, 2015 at 2:44 PM, Artem Belevich via cfe-commits <span dir="ltr"><<a href="mailto:cfe-commits@lists.llvm.org" target="_blank">cfe-commits@lists.llvm.org</a>></span> wrote:<br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">tra created this revision.<br>
tra added reviewers: rsmith, jingyue, jpienaar.<br>
tra added a subscriber: cfe-commits.<br>
<br>
C++ emits vtables for classes that have key function present in the<br>
current TU. While we compile CUDA the fact that key function was found<br>
in this TU does not mean that we are going to generate code for it. E.g.<br>
vtable for a class with host-only methods should not be generated on<br>
device side, because we are not going to generate any code for the<br>
host-only methods during device-side compilation.<br>
<br>
During CUDA compilation this patch checks virtual methods' target attributes and<br>
returns key function only if all virtual methods in the class are suitable for the current<br>
compilation mode.<br>
<br>
<br>
<a href="http://reviews.llvm.org/D15309" rel="noreferrer" target="_blank">http://reviews.llvm.org/D15309</a><br>
<br>
Files:<br>
lib/AST/RecordLayoutBuilder.cpp<br>
test/CodeGenCUDA/<a href="http://device-vtable.cu" rel="noreferrer" target="_blank">device-vtable.cu</a><br>
<br>
Index: test/CodeGenCUDA/<a href="http://device-vtable.cu" rel="noreferrer" target="_blank">device-vtable.cu</a><br>
===================================================================<br>
--- /dev/null<br>
+++ test/CodeGenCUDA/<a href="http://device-vtable.cu" rel="noreferrer" target="_blank">device-vtable.cu</a><br>
@@ -0,0 +1,55 @@<br>
+// REQUIRES: x86-registered-target<br>
+// REQUIRES: nvptx-registered-target<br>
+<br>
+// Make sure we don't emit vtables for classes with methods that have<br>
+// inappropriate target attributes. Currently it's mostly needed in<br>
+// order to avoid emitting vtables for host-only classes on device<br>
+// side where we can't codegen them.<br>
+<br>
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -o - %s \<br>
+// RUN: | FileCheck %s -check-prefix=CHECK-HOST<br>
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -o - %s \<br>
+// RUN: | FileCheck %s -check-prefix=CHECK-DEVICE<br>
+<br>
+#include "Inputs/cuda.h"<br>
+<br>
+class H {<br>
+ public:<br>
+ virtual void method();<br>
+};<br>
+//CHECK-HOST: @_ZTV1H =<br>
+//CHECK-HOST-SAME: @_ZN1H6methodEv<br>
+//CHECK-DEVICE-NOT: @_ZTV1H =<br>
+<br>
+class D {<br>
+ public:<br>
+ __device__ virtual void method();<br>
+};<br>
+<br>
+//CHECK-DEVICE: @_ZTV1D<br>
+//CHECK-DEVICE-SAME: @_ZN1D6methodEv<br>
+//CHECK-HOST-NOT: @_ZTV1D<br>
+<br>
+// This is the case with mixed host and device virtual methods. It's<br>
+// impossible to emit a valid vtable in that case because only host or<br>
+// only device methods would be available during host or device<br>
+// compilation. For now we'll not emit such vtable at all.<br>
+class HD {<br>
+ public:<br>
+ virtual void h_method();<br>
+ __device__ virtual void d_method();<br>
+};<br>
+//CHECK-BOTH-NOT: @_ZTV2HD<br>
+<br>
+void H::method() {}<br>
+//CHECK-HOST: define void @_ZN1H6methodEv<br>
+<br>
+void __device__ D::method() {}<br>
+//CHECK-DEVICE: define void @_ZN1D6methodEv<br>
+<br>
+void __device__ HD::d_method() {}<br>
+// CHECK-DEVICE: define void @_ZN2HD8d_methodEv<br>
+// CHECK-HOST-NOT: define void @_ZN2HD8d_methodEv<br>
+void HD::h_method() {}<br>
+// CHECK-HOST: define void @_ZN2HD8h_methodEv<br>
+// CHECK-DEVICE-NOT: define void @_ZN2HD8h_methodEv<br>
Index: lib/AST/RecordLayoutBuilder.cpp<br>
===================================================================<br>
--- lib/AST/RecordLayoutBuilder.cpp<br>
+++ lib/AST/RecordLayoutBuilder.cpp<br>
@@ -1996,6 +1996,16 @@<br>
bool allowInlineFunctions =<br>
Context.getTargetInfo().getCXXABI().canKeyFunctionBeInline();<br>
<br>
+ if (Context.getLangOpts().CUDA) {<br>
+ const bool IsDevice = Context.getLangOpts().CUDAIsDevice;<br></blockquote><div><br></div><div>Guess you intended to use this local variable, but didn't? ^</div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
+ for (const CXXMethodDecl *MD : RD->methods())<br>
+ if (Context.getLangOpts().CUDAIsDevice && !MD->hasAttr<CUDADeviceAttr>())<br>
+ return nullptr;<br>
+ else if (!Context.getLangOpts().CUDAIsDevice &&<br></blockquote><div><br></div><div>Drop the else after return (just "if/return/if/return")<br>(and/or you could roll both conditions into one if test - potentially even using "std::any_of" (or llvm::any_of))</div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
+ !MD->hasAttr<CUDAHostAttr>() && MD->hasAttr<CUDADeviceAttr>())<br>
+ return nullptr;<br>
+ }<br>
+<br>
for (const CXXMethodDecl *MD : RD->methods()) {<br>
if (!MD->isVirtual())<br>
continue;<br>
<br>
<br>
<br>_______________________________________________<br>
cfe-commits mailing list<br>
<a href="mailto:cfe-commits@lists.llvm.org" target="_blank">cfe-commits@lists.llvm.org</a><br>
<a href="http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits" rel="noreferrer" target="_blank">http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits</a><br>
<br></blockquote></div><br></div></div>