[PATCH] D12241: [CUDA] Change initializer for CUDA device code based on CUDA documentation.

Artem Belevich via llvm-commits llvm-commits at lists.llvm.org
Fri Aug 21 10:23:27 PDT 2015


tra added a comment.

+1 to jingyue's comments.


================
Comment at: lib/CodeGen/CodeGenModule.cpp:1995
@@ +1994,3 @@
+  // of their declaration."
+  if (LangOpts.CUDA && LangOpts.CUDAIsDevice
+      && D->hasAttr<CUDASharedAttr>()) {
----------------
please use getLangOpts()

================
Comment at: lib/CodeGen/CodeGenModule.cpp:2086-2096
@@ -2078,1 +2085,13 @@
 
+  // CUDA B.2.1 "The __device__ qualifier declares a variable that resides on
+  // the device. [...]
+  // Is accessible from all the threads within the grid and from the host
+  // through the runtime library (cudaGetSymbolAddress() / cudaGetSymbolSize()
+  // / cudaMemcpyToSymbol() / cudaMemcpyFromSymbol()).
+  //
+  // CUDA B.2.2 "The __constant__ qualifier, optionally used together with
+  // __device__, declares a variable that: [...]
+  // Is accessible from all the threads within the grid and from the host
+  // through the runtime library (cudaGetSymbolAddress() / cudaGetSymbolSize()
+  // / cudaMemcpyToSymbol() / cudaMemcpyFromSymbol()).
+  if (GV && LangOpts.CUDA && LangOpts.CUDAIsDevice &&
----------------
Cosmetic nit: no closing quote marks in both citations.

You may want to leave only one copy if "is accessible..." as it applies to both B.2.1 and B.2.2.


http://reviews.llvm.org/D12241





More information about the llvm-commits mailing list