Add shared mem-space for SLM allocs #397
Conversation
Signed-off-by: dchigarev <dmitry.chigarev@intel.com>
Signed-off-by: dchigarev <dmitry.chigarev@intel.com>
| while (parentOp) { | ||
| // Check if the current parent is a gpu.launch operation | ||
| if (llvm::isa<mlir::gpu::LaunchOp>(parentOp)) { | ||
| return true; | ||
| } | ||
| // Move to the parent operation | ||
| parentOp = parentOp->getParentOp(); | ||
| } |
There was a problem hiding this comment.
You're probably looking for getParentOfType here. Another option would be to just target the pass at gpu module instead. The latter is probably also useful if you want to check for the total allocated memory size without additional analysis.
There was a problem hiding this comment.
Another option would be to just target the pass at gpu module instead
Targeting the pass to a gpu module would be nice, however it won't work with the existing pipeline. There are some passes that must be run after allocs-to-slm but before conversion to the gpu module (linalg-to-xegpu, linalg-to-loops, insert-gpu-allocs).
Some of the passes (linalg-to-xegpu and linalg-to-loops) might be adjusted by retargeting them from funcOp to gpuFuncOp, but insert-gpu-allocs won't work that way, as its logic expects to be applied to a func.func containing a gpu.launch block (and not a GPU module).
There was a problem hiding this comment.
So only replaced the while-loop with getParentOfType for now
There was a problem hiding this comment.
Sorry, I was thinking about something else while typing. I meant the launch gpu op, not gpu module. Anyway, the change is trivial.
Signed-off-by: dchigarev <dmitry.chigarev@intel.com>
Signed-off-by: dchigarev <dmitry.chigarev@intel.com>
3 participants
Fixes #393
The new pass goes through each
memref.alloc()inside of agpu.launch block {}and sets shared mem-space attribute (int value of 3) to the resulted memref (it ignores all memrefs that already have explicitly specified mem-space).