[Mlir-commits] [mlir] [mlir][gpu] Introduce `gpu.dynamic_shared_memory` Op (PR #71546)

[Guray Ozen]

================
@@ -554,6 +555,95 @@ static IntegerAttr wrapNumericMemorySpace(MLIRContext *ctx, unsigned space) {
   return IntegerAttr::get(IntegerType::get(ctx, 64), space);
 }
 
+/// Generates a symbol with 0-sized array type for dynamic shared memory usage,
+/// or uses existing symbol.
+LLVM::GlobalOp
+getDynamicSharedMemorySymbol(ConversionPatternRewriter &rewriter,
+                             gpu::DynamicSharedMemoryOp op,
+                             const LLVMTypeConverter *typeConverter,
+                             MemRefType memrefType, unsigned alignmentBit) {
+  std::optional<LLVM::GlobalOp> existingGlobalOp;
+
+  LLVM::LLVMFuncOp funcOp = op->getParentOfType<LLVM::LLVMFuncOp>();
+  assert(funcOp && "cannot find llvm.func op");
+
+  gpu::GPUModuleOp moduleOp = funcOp->getParentOfType<gpu::GPUModuleOp>();
+  assert(moduleOp && "cannot find gpu.module op");
+
+  // Use already generated global op if it exists
+  int index = 0;
+  std::string prefix = llvm::formatv("__shmem_{0}", funcOp.getSymName());
+  moduleOp->walk([&](LLVM::GlobalOp globalOp) {
+    if (auto arrayType = dyn_cast<LLVM::LLVMArrayType>(globalOp.getType())) {
+      if (arrayType.getNumElements() == 0) {
+        existingGlobalOp = globalOp;
+        return WalkResult::interrupt();
+      }
+    }
+    if (globalOp.getSymName().startswith(prefix))
+      index++;
----------------
grypp wrote:

Agreed using SymbolTable! This was a workaround for name clashing. While building the symbol table with the GPUModule, I am getting the following errors. I will look deeper into this tomorrow, but if you can spot any obvious issues, please let me know 


```
gpu::GPUModuleOp moduleOp = funcOp->getParentOfType<gpu::GPUModuleOp>();
SymbolTable table(moduleOp);
```

```
mlir-opt: /usr/local/google/home/gurayozen/work/llvm-project/mlir/lib/IR/SymbolTable.cpp:137: mlir::SymbolTable::SymbolTable(mlir::Operation *): Assertion `inserted.second && "expected region to contain uniquely named symbol operations"' failed.
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0.      Program arguments: mlir-opt dynamic-shared-memory.mlir -convert-gpu-to-nvvm -cse -canonicalize -o asd
----
#10 0x0000556d3e0c3e43 mlir::SymbolTable::SymbolTable(mlir::Operation*) /usr/local/google/home/gurayozen/work/llvm-project/mlir/lib/IR/SymbolTable.cpp:129:3
#11 0x0000556d3cbedf15 getDynamicSharedMemorySymbol(mlir::ConversionPatternRewriter&, mlir::gpu::DynamicSharedMemoryOp, mlir::LLVMTypeConverter const*, mlir::MemRefType, unsigned int) /usr/local/google/home/gurayozen/work/llvm-project/mlir/lib/Conversion/GPUCommon/GPUOpsLowering.cpp:577:3
#12 0x0000556d3cbee362 mlir::GPUDynamicSharedMemoryOpLowering::matchAndRewrite(mlir::gpu::DynamicSharedMemoryOp, mlir::gpu::DynamicSharedMemoryOpAdaptor, mlir::ConversionPatternRewriter&) const /usr/local/google/home/gurayozen/work/llvm-project/mlir/lib/Conversion/GPUCommon/GPUOpsLowering.cpp:619:28
---
```

https://github.com/llvm/llvm-project/pull/71546