[PATCH] D17779: [CUDA] Emit host-side 'shadows' for device-side global variables

[Justin Lebar via cfe-commits]

jlebar added inline comments.

================
Comment at: lib/CodeGen/CGCUDANV.cpp:168
@@ -163,1 +167,3 @@
+/// of global scope device-side variables generated in this module
+/// with the CUDA runtime.
 /// \code
----------------
This is kind of hard to parse.  How about rephrasing to something like:

Creates a function that sets up state on the host side for CUDA objects that have a presence on both the host and device sides.  Specifically, registers the host side of kernel functions and __device__ global variables with the CUDA runtime.

================
Comment at: lib/CodeGen/CGCUDANV.cpp:213
@@ +212,3 @@
+  // void __cudaRegisterVar(void **, char *, char *, const char *,
+  //                        int, int, int, int)
+  std::vector<llvm::Type *> RegisterVarParams = {
----------------
Can we say what these args mean?

================
Comment at: lib/CodeGen/CGCUDANV.cpp:224
@@ +223,3 @@
+    llvm::Constant *VarName = makeConstantString(Var->getName());
+    llvm::Value *args[] = {
+        &GpuBinaryHandlePtr, Builder.CreateBitCast(Var, VoidPtrTy), VarName,
----------------
Nit: s/args/Args/?

================
Comment at: lib/CodeGen/CGCUDANV.cpp:228
@@ +227,3 @@
+        llvm::ConstantInt::get(IntTy, CGM.getDataLayout().getTypeAllocSize(
+                                          Var->getValueType())), // sizeof(var)
+        llvm::ConstantInt::get(IntTy, (Flags & DevVarConst) ? 1 : 0),
----------------
Nit: Maybe pull this expression out as a separate var?  Then the comment isn't needed (would be nice, because at the moment it's ambiguous exactly what "sizeof(var)" refers to.

================
Comment at: lib/CodeGen/CodeGenModule.cpp:1532
@@ +1531,3 @@
+      // We need to emit host-side 'shadows' for all global
+      // device-side variables because CUDA runtime API needs their
+      // size and host-side address in order to provide access to
----------------
s/CUDA runtime API/the CUDA runtime/ (not really a requirement of the API, I think?)

================
Comment at: lib/CodeGen/CodeGenModule.cpp:1575
@@ +1574,3 @@
+      // definition, because we still need to define host-side shadow
+      // for it.
+    } else if (VD->isThisDeclarationADefinition() != VarDecl::Definition &&
----------------
Kind of an odd way of writing this control flow?  Could we phrase it more idiomatically as

  MustEmitForCUDA = !VD->hasDefinition() && ...;
  if (!MustEmitForCUDA && ...) return;

================
Comment at: lib/CodeGen/CodeGenModule.cpp:2477
@@ +2476,3 @@
+      if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()) {
+        Linkage = llvm::GlobalValue::InternalLinkage;
+
----------------
Is it worth explaining why the shadows get internal linkage?

================
Comment at: lib/CodeGen/CodeGenModule.cpp:2480
@@ +2479,3 @@
+        // Shadow variables and their properties must be registered
+        // with CUDA runtime.
+        unsigned Flags = 0;
----------------
with the CUDA runtime

================
Comment at: lib/CodeGen/CodeGenModule.cpp:2483
@@ +2482,3 @@
+        if (!D->hasDefinition())
+          Flags |= CGCUDARuntime::DevVarExt;
+        if (D->hasAttr<CUDAConstantAttr>())
----------------
Now that I see them in context, I think these flags would be a lot easier to handle if they employed less abbreviation.  "ExternalDeviceVar", "ConstDeviceVar"?

================
Comment at: test/CodeGenCUDA/device-stub.cu:14
@@ +13,3 @@
+
+// Make sure host globals don't get internalized..
+// CHECK-DAG: @host_var = global i32
----------------
Not sure if this is a typo or if you mean "...".

================
Comment at: test/CodeGenCUDA/device-stub.cu:17
@@ +16,3 @@
+int host_var;
+// .. and that extern vars remain external.
+// CHECK-DAG: @ext_host_var = external global i32
----------------
Here you do seem to mean "..."


http://reviews.llvm.org/D17779