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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits