[PATCH] D91516: [AMDGPU][WIP] Lower Function Local LDS Variables.

[Mahesha S via Phabricator via llvm-commits]

hsmhsm marked 22 inline comments as done.
hsmhsm added inline comments.


================
Comment at: llvm/lib/Target/AMDGPU/AMDGPULowerFunctionLocalLDS.cpp:111
+    for (auto *LDS : LDSGlobals)
+      if (LDSToFunction.find(LDS) == LDSToFunction.end())
+        ToBeRemovedLDSList.insert(LDS);
----------------
arsenm wrote:
> .contains
The data structures, ValueMap<>,  SmallPtrSet<> do not have member function -  `.contains()`. W.r.t std::set<>, this member function is supported in C++20.


================
Comment at: llvm/lib/Target/AMDGPU/AMDGPULowerFunctionLocalLDS.cpp:376-379
+  // Push the `CallGraphNode` associated with all the callees of the kernel`K`
+  // into into `CGNodeStack`, and the corresponding call sites into
+  // `CallBaseStack`.
+  pushCallGraphNodes(KernCGNode, CGNodeStack, CallBaseStack);
----------------
arsenm wrote:
> I think you're overcomplicating the CallGraph usage by ignoring most of what it gives you. You should be able to just iterate directly through the CallGraph to get functions reachable from the parent
As far as I understand it, llvm `CallGraph` infrastructure does not provide any facility as such.  Implementer needs to explicitly iterate the callees of the caller.


================
Comment at: llvm/lib/Target/AMDGPU/AMDGPULowerFunctionLocalLDS.cpp:410-411
+
+// Create a reverse map from function to LDS set which maps a given function F
+// to a set of LDS which are defined within F.
+void LowerFunctionLocalLDSImpl::createFunctionToLDSMap() {
----------------
arsenm wrote:
> This concept doesn't quite work for the IR. The same global can appear in multiple functions
My understanding is that - scope of the shared variable is function/statement block scope. It is not available to access outside this scope. It is just that we implement it as global, just like how the local static variables are implemented in C/C++?

Can you give an example of the use-case that you are claiming?


================
Comment at: llvm/lib/Target/AMDGPU/AMDGPULowerFunctionLocalLDS.cpp:502
+    if (GV.getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS &&
+        !GV.hasExternalLinkage())
+      LDSGlobals.insert(&GV);
----------------
arsenm wrote:
> Probably should skip declarations. Also not sure about the linkage check
The linkage test is required to ignore the `dynamic shared variables` like the one defined as `extern __shared__ int dy_sm[];`  where size of `dy_sm` is not known at compile time, but is passed as one of the kernel execution configuration parameters at run time.


================
Comment at: llvm/lib/Target/AMDGPU/AMDGPULowerFunctionLocalLDS.cpp:561-571
+// Entry-point function.
+bool LowerFunctionLocalLDSImpl::lower() {
+  // Build necessary data structures which will be later used in the lowering
+  // process.
+  if (!initialize())
+    return false;
+
----------------
arsenm wrote:
> Don't understand the point of this stub function
This is a driver function, it looks like a stub now, since implementation is not complete yet. Once this patch is accepted, next step is to (1) define kernel specific LDS layouts (2) create 2D offset table and (3) add new implicit argument.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D91516/new/

https://reviews.llvm.org/D91516