[llvm] [MLIR] Enabling Intel GPU Integration. (PR #65539)

Ronan Keryell via llvm-commits llvm-commits at lists.llvm.org
Thu Sep 7 18:42:00 PDT 2023


================
@@ -811,8 +812,13 @@ LogicalResult ConvertAllocOpToGpuRuntimeCallPattern::matchAndRewrite(
   // descriptor.
   Type elementPtrType = this->getElementPtrType(memRefType);
   auto stream = adaptor.getAsyncDependencies().front();
+
+  auto isHostShared = rewriter.create<mlir::LLVM::ConstantOp>(
+      loc, llvmInt64Type, rewriter.getI64IntegerAttr(isShared));
+
   Value allocatedPtr =
-      allocCallBuilder.create(loc, rewriter, {sizeBytes, stream}).getResult();
+      allocCallBuilder.create(loc, rewriter, {sizeBytes, stream, isHostShared})
+          .getResult();
----------------
keryell wrote:

Technically, SYCL provides a more abstract memory management with `sycl::buffer` and `sycl::accessor` defining an implicit asynchronous task graph. The allocation details are left to the implementation, asynchronous or synchronous allocation is left to the implementers.
Here the lower-level synchronous USM memory management API of SYCL is used instead, similar to CUDA/HIP memory management.
So, should the `async` allocation in the example be synchronous instead?

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


More information about the llvm-commits mailing list