[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