[PATCH] D38978: [OpenMP] Enable the lowering of implicitly shared variables in OpenMP GPU-offloaded target regions to the GPU shared memory

[Artem Belevich via Phabricator via cfe-commits]

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