[PATCH] D147172: [OpenMP][Flang][MLIR] Add lowering of TargetOp for host codegen

[Johannes Doerfert via Phabricator via llvm-commits]

jdoerfert added a comment.

TL;DR, we should try to converge to one impl. in OpenMPIRBuilder but that might take a while and we should not force it where it doesn't make sense (yet).



================
Comment at: llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h:1829
+  using TargetBodyGenCallbackTy = function_ref<InsertPointTy(
+      InsertPointTy AllocaIP, InsertPointTy CodeGenIP)>;
+  /// Generator for '#omp target'
----------------
Nit: newline and docs above; copy&paste error below; it's not the default number but clause value;


================
Comment at: llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp:4166
+
+  OpenMPIRBuilder::FunctionGenCallback &&generateOutlinedFunction =
+      [&Builder, &Inputs, &CBFunc](StringRef EntryFnName) {
----------------



================
Comment at: llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp:4180
+  // TODO: Add kernel launch call when device codegen is supported.
+  Builder.CreateCall(OutlinedFn, Args);
+}
----------------
FWIW, you should already (be able to) emit the kernel launch call and fallback handling.
The latter should kick in as the symbol is neither registered nor available on the device.


================
Comment at: mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp:1554
+
+  assert(fileLoc && "No file found from location");
+  StringRef fileName = fileLoc.getFilename().getValue();
----------------
jsjodin wrote:
> skatrak wrote:
> > Is it generally expected for the compilation to fail if the input is given via stdin? It looks like clang currently does implement this behavior but it's not communicated to the user:
> > ```
> > $ echo -e "int main(){return 0;}" | clang -cc1 -fopenmp - -emit-llvm -o -
> > ; ModuleID = '-'
> > source_filename = "-"
> > target datalayout = ...
> > ...
> > 
> > $ echo -e "int main(){\n#pragma omp target\n{} return 0;}" | clang -cc1 -fopenmp - -emit-llvm -o -
> > fatal error: cannot open file '<stdin>': No such file or directory
> > 
> > ```
> > 
> > Currently flang-new seems to support compiling from stdin. Or at least that's the case for the simplest program:
> > ```
> > $ echo "end program" | flang-new -fc1 -x f95-cpp-input - -emit-mlir -o -
> > module attributes {dlti.dl_spec ...
> > ...
> > ```
> Yes, the assumption in clang is that there is a file, which is needed to generate a unique name for outlined target region functions. The code above would make flang-new fail the same way if target regions are used.
@skatrak You can file a bug. We might be able to handle that fine, but it has not come up yet.


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

https://reviews.llvm.org/D147172