tra added a comment.
Please add tests for the cases where such local->shaed conversion should and
should not happen.
I would appreciate if you could add details on what exactly your passes are
supposed to move to shared memory.
Considering that device-side code tends to be heavily inlined, it may be
prudent to add an option to control the total size of shared memory we allow to
be used for this purpose.
In case your passes are not executed (or didn't move anything to shared
memory), is there any impact on the generated PTX. I.e. can ptxas successfully
optimize unused shared memory away?
If the code intentionally wants to allocate something in local memory, would
the allocation ever be moved to shared memory by your pass? If so, how would I
prevent that?
================
Comment at: lib/Target/NVPTX/NVPTXAsmPrinter.cpp:1749
+ O << "\t.reg .b32 \t%SHSP;\n";
+ O << "\t.reg .b32 \t%SHSPL;\n";
+ }
----------------
Nit: the name should end with `S` as the L in `SPL` was for 'local' address
space. which then gets converted to generic AS. In your case it will be in
shared space, hence S would be more appropriate.
================
Comment at: lib/Target/NVPTX/NVPTXAssignValidGlobalNames.cpp:68
+void NVPTXAssignValidGlobalNames::generateCleanName(Value &V) {
+ std::string ValidName;
----------------
The name cleanup changes in this file should probably be committed by
themselves as they have nothing to do with the rest of the patch.
================
Comment at: lib/Target/NVPTX/NVPTXFunctionDataSharing.cpp:9
+//===----------------------------------------------------------------------===//
+//
+//
----------------
Please add details about what the pass is supposed to do.
Repository:
rL LLVM
https://reviews.llvm.org/D38978
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
