jlebar accepted this revision. jlebar added a comment. This revision is now accepted and ready to land.
LGTM, mostly nits. ================ Comment at: clang/include/clang/Sema/Sema.h:10316 + /// Returns the name of the launch configuration function. + std::string getCudaConfigureFuncName() const; ---------------- Could we be a little less vague, what exactly is the launch-configuration function? (Could be as simple as adding `e.g. cudaFooBar()`.) ================ Comment at: clang/lib/CodeGen/CGCUDANV.cpp:201 -void CGNVCUDARuntime::emitDeviceStubBody(CodeGenFunction &CGF, - FunctionArgList &Args) { +// CUDA 9.0+ uses new way to launch kernels. Parameters are packed in local +// array and kernels are launched using cudaLaunchKernel(). ---------------- nit `in a local array` ================ Comment at: clang/lib/CodeGen/CGCUDANV.cpp:212 + VoidPtrTy, CharUnits::fromQuantity(16), "kernel_args", + llvm::ConstantInt::get(SizeTy, std::max(1UL, Args.size()))); + // Store pointers to the arguments in a locally allocated launch_args. ---------------- Nit, s/`1UL`/`uint64{1}`/ or size_t, whatever this function takes. As-is we're baking in the assumption that unsigned long is the same as the type returned by Args.size(), which isn't necessarily true. As an alternative, you could do `std::max<size_t>(1, Args.size())` or whatever the appropriate type is. ================ Comment at: clang/lib/CodeGen/CGCUDANV.cpp:239 + CGM.Error(CGF.CurFuncDecl->getLocation(), + "Can't find declaration for cudaLaunchKernel()"); // FIXME. + return; ---------------- Unfixed FIXME? ================ Comment at: clang/lib/CodeGen/CGCUDANV.cpp:260 + /*isVarArg=*/false), + "__cudaPopCallConfiguration"); + ---------------- I see lots of references to `__cudaPushCallConfiguration`, but this is the only reference I see to `__cudaPopCallConfiguration`. Is this a typo? Also are we supposed to emit matching push and pop function calls? Kind of weird to do one without the other... ================ Comment at: clang/lib/CodeGen/CGCUDANV.cpp:266 + // Emit the call to cudaLaunch + + llvm::Value *Kernel = CGF.Builder.CreatePointerCast(CGF.CurFn, VoidPtrTy); ---------------- Whitespace nit, maybe move this whitespace line before the comment? ================ Comment at: clang/lib/Headers/__clang_cuda_runtime_wrapper.h:429 +// CUDA runtime uses undocumented function to access kernel launch +// configuration. We need to provide our own declaration for it here. ---------------- s/undocumented function/this undocumented function/? CHANGES SINCE LAST ACTION https://reviews.llvm.org/D57488/new/ https://reviews.llvm.org/D57488 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits