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

[jsjodin]

Summary

This patch adds lowering of TargetOps for the host. The lowering outlines the
target region function and uses the OpenMPIRBuilder support functions to emit
the function and call. Code generation for offloading will be done in
later patches. The tests only use scalars for now to reduce the size of the initial
patches for offloading. Runtime tests will be added in a separate patch since 'lit'
is currently not set up to handle Fortran runtime tests.

[kiranchandramohan]

Please split this patch into three patches (or two).

1. Containing just the OpenMPIRBuilder changes.
2. Containing the translation from OPenMP + LLVM Dialect to LLVM IR.
3. The conversion from Fortran to FIR + OpenMP and the test from FIR + OpenMP to LLVM + OpenMP.

See comments inline. We would like to have the translation layer thin and have most of the code in the OpenMP IRBuilder. Please refer to the translation of other OpenMP Ops as detailed in the inline comment. If there is anything special to this Op that prevents it then please explain.

flang/test/Lower/OpenMP/target_region_to_llvmir.f90

1 ↗

(On Diff #509427)

We generally do not test lowering to LLVM IR. It is usually done in multiple steps:

1. Test for lowering from Fortran+OpenMP source to FIR + OpenMP dialects in flang/test/Lower/OpenMP/.
2. Test conversion of FIR + OpenMP dialects to LLVM + OpenMP dialects in lang/test/Fir/convert-to-llvm-openmp-and-fir.fir.
3. Test translation of OpenMP + LLVM Dialect to LLVM IR in mlir/test/Target/LLVMIR/openmp-*.mlir.

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

1664

We would generally like to keep the Translation code thin and do the heavy-lifting in the OpenMP IRBuilder. I think this should be rewritten with more code sunk into the OpenMP IRBuilder. Can this be converted to other functions like convertOmpSingle or convertOmpMaster or convertOmpTaskOp or convertOmpParallel?

mlir/lib/Target/LLVMIR/ModuleTranslation.cpp

1349–1351

Nit: Braces not required.

1353

Nit: What are the three false options?

[jsjodin]

flang/test/Lower/OpenMP/target_region_to_llvmir.f90

1 ↗

(On Diff #509427)

Yes, I will split up the patch into 3 separate ones. What is the expectation on the various clauses? Can these be implemented in separate sets of patches? The patches going from OMP+LLVM Dialect -> LLVM IR could be very large if we do everything at once.

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

1664

I agree we want to put as much as possible in the OpenMPIRBuilder, but unless both clang and MLIR can use it, it shouldn't be in the OpenMPIRBuilder, because it would expose two different (possibly incompatible) implementations of the same thing to either clang or mlir, which in my opinion is undesirable. When we did the migration of code from clang to OpenMPIRBuilder we tried to migrate as much as possible, the problem with generating code for the target regions specifically is that there are data structures that need to be maintained which all use clang data types like mapping device pointers to values etc. There are also issues with clauses affecting the logic for how the code should be generated, these two things combined are not easy to get around. Parameterizing the code and using callback functions (which OpenMPIRBuilder does today) didn't seem feasible, because the common code would be trivial, and the parameterization would just make things less readable/maintainable. Another idea would be to use templates to be able to ignore the specific types, which might be a way forward, but unless we know what the MLIR implementation looks like it is hard to know if it would work. Another potential alternative would be to have some kind of interface for dealing with values/statements/clauses which the OpenMPIRBuilder would define and clang/mlir implement. But after some discussions with @jdoerfert we concluded that this would probably not be a good option (too complicated basically). Hopefully something will emerge while we are working on the MLIR side. If sharing code through the OpenMPIRBuilder is not a good possible, then it might be better to keep it in MLIR,. What are your thoughts about this?

mlir/lib/Target/LLVMIR/ModuleTranslation.cpp

1353

Nit: What are the three false options?

Because we don't have the flags/dierctives implemented in the frontend yet.

[kiranchandramohan]

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

1664

Thanks for the detailed explanation here.

Unless there are issues for Flang/MLIR, I would still recommend moving the code to the OpenMPIRBuilder for a few reasons:

1. Keep the translation of the target directive consistent with the others. Outlining and insertion of runtime calls are currently performed by the OpenMPIRBuilder. Having it in translation will confuse future contributors. Also, in the present state, it looks like the translation for the target operation can follow the same pattern as the operations.
2. Having the Flang/MLIR version in the OpenMPIRBuilder is beneficial for any out-of-tree projects as well. There are already other constructs where Clang is not using the OpenMPIRBuilder. And having two interfaces that are well-documented is better than not having any.
3. Keep the code in translation as simple as possible. This is what we promised and agreed on with the MLIR folks.

mlir/lib/Target/LLVMIR/ModuleTranslation.cpp

1353

I mean't, please add a comment for each of the options so that someone who reads the code knows what they are.

[skatrak]

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

1554

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 ...
...

[kiranchandramohan]

flang/test/Lower/OpenMP/target_region_to_llvmir.f90

1 ↗

(On Diff #509427)

Yes, you can split the clauses into a series of patches.

[jsjodin]

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

1664

Unless there are issues for Flang/MLIR, I would still recommend moving the code to the OpenMPIRBuilder for a few reasons:

1. Keep the translation of the target directive consistent with the others. Outlining and insertion of runtime calls are currently performed by the OpenMPIRBuilder. Having it in translation will confuse future contributors. Also, in the present state, it looks like the translation for the target operation can follow the same pattern as the operations.

Yes, there should be a way to use the same pattern, I can work on creating a convertOmpTarget function in the OpenMPIRBuilder that will do this.

2. Having the Flang/MLIR version in the OpenMPIRBuilder is beneficial for any out-of-tree projects as well. There are already other constructs where Clang is not using the OpenMPIRBuilder. And having two interfaces that are well-documented is better than not having any.

My initial thinking was to keep the code structure similar between MLIR and Clang similar to make it easier to have a shared OpenMPIRBuilder interface. Given 3. there isn't much choice than to add new functions in OpenMPIRBuilder to minimize the code in MLIR.

3. Keep the code in translation as simple as possible. This is what we promised and agreed on with the MLIR folks.

Good to know!

[jsjodin]

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

1554

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.

[jsjodin]

@kiranchandramohan I moved the outlining to the OpenMPIRBuilder, is this more in line of what you were thinking? There is a patch up by @TIFitis for Fortran->Fir lowering https://reviews.llvm.org/D147339. I will prune this patch to only have the OMP+LLVM Dialect -> LLVMIR portions and appropriate tests.

[jsjodin]

Fix formatting and remove debug code.

[TIFitis]

For IRBuilder lowering unit tests need to be added here: llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp
And for MLIR Translation tests need to be added to: mlir/test/Target/LLVMIR/omptarget-llvm.mlir

[kiranchandramohan]

In D147172#4237283, @jsjodin wrote:

@kiranchandramohan I moved the outlining to the OpenMPIRBuilder, is this more in line of what you were thinking? There is a patch up by @TIFitis for Fortran->Fir lowering https://reviews.llvm.org/D147339. I will prune this patch to only have the OMP+LLVM Dialect -> LLVMIR portions and appropriate tests.

Thanks @jsjodin for the changes. Yes, moving the code to the OpenMPIRBuilder looks fine. I have not looked at the patch in detail. But I see that you have not used the OpenMPIRBuilder's outlining facility (OutlineInfo(s) and the PostOutlineCB). Are there any issues in using these here?

[jsjodin]

In D147172#4240342, @kiranchandramohan wrote:

In D147172#4237283, @jsjodin wrote:

@kiranchandramohan I moved the outlining to the OpenMPIRBuilder, is this more in line of what you were thinking? There is a patch up by @TIFitis for Fortran->Fir lowering https://reviews.llvm.org/D147339. I will prune this patch to only have the OMP+LLVM Dialect -> LLVMIR portions and appropriate tests.

Thanks @jsjodin for the changes. Yes, moving the code to the OpenMPIRBuilder looks fine. I have not looked at the patch in detail. But I see that you have not used the OpenMPIRBuilder's outlining facility (OutlineInfo(s) and the PostOutlineCB). Are there any issues in using these here?

I have reworked the patch a bit to simplify things and added tests. The reason why the OutlineInfos etc are not used for the target regions is probably because the outlining is more complicated. With offloading we have to generate an if-then-else that tries to offload via the openmp runtime and if the runtime fails to launch the kernel the fallback is executed instead.

[jsjodin]

Simplify the code by only passing the inputs, since they will be the same values in the outlined code. Add tests and improve comments.

[jsjodin]

In D147172#4244736, @jsjodin wrote:

In D147172#4240342, @kiranchandramohan wrote:

In D147172#4237283, @jsjodin wrote:

@kiranchandramohan I moved the outlining to the OpenMPIRBuilder, is this more in line of what you were thinking? There is a patch up by @TIFitis for Fortran->Fir lowering https://reviews.llvm.org/D147339. I will prune this patch to only have the OMP+LLVM Dialect -> LLVMIR portions and appropriate tests.

Thanks @jsjodin for the changes. Yes, moving the code to the OpenMPIRBuilder looks fine. I have not looked at the patch in detail. But I see that you have not used the OpenMPIRBuilder's outlining facility (OutlineInfo(s) and the PostOutlineCB). Are there any issues in using these here?

I have reworked the patch a bit to simplify things and added tests. The reason why the OutlineInfos etc are not used for the target regions is probably because the outlining is more complicated. With offloading we have to generate an if-then-else that tries to offload via the openmp runtime and if the runtime fails to launch the kernel the fallback is executed instead.

Also since the outlining is shared with the device codegen, that is also a reason why the OutlineInfos etc is not used.

[TIFitis]

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

1575–1576

Can you add checks for clauses that are not yet supported?

mlir/test/Target/LLVMIR/omptarget-llvm.mlir

178

I think the module can be omitted from the test.

[kiranchandramohan]

In D147172#4245879, @jsjodin wrote:

In D147172#4244736, @jsjodin wrote:

In D147172#4240342, @kiranchandramohan wrote:

In D147172#4237283, @jsjodin wrote:

@kiranchandramohan I moved the outlining to the OpenMPIRBuilder, is this more in line of what you were thinking? There is a patch up by @TIFitis for Fortran->Fir lowering https://reviews.llvm.org/D147339. I will prune this patch to only have the OMP+LLVM Dialect -> LLVMIR portions and appropriate tests.

Thanks @jsjodin for the changes. Yes, moving the code to the OpenMPIRBuilder looks fine. I have not looked at the patch in detail. But I see that you have not used the OpenMPIRBuilder's outlining facility (OutlineInfo(s) and the PostOutlineCB). Are there any issues in using these here?

I have reworked the patch a bit to simplify things and added tests. The reason why the OutlineInfos etc are not used for the target regions is probably because the outlining is more complicated. With offloading we have to generate an if-then-else that tries to offload via the openmp runtime and if the runtime fails to launch the kernel the fallback is executed instead.

Also since the outlining is shared with the device codegen, that is also a reason why the OutlineInfos etc is not used.

We have handling for the if clause in both task and parallel. Will the handling in Target be substantially different from these?
Task : D130615
Parallel : D138495 (The latest revision uses a runtime extension, so that we do not have to generate IR for both sequential and parallel version).

[jsjodin]

In D147172#4246212, @kiranchandramohan wrote:

In D147172#4245879, @jsjodin wrote:

In D147172#4244736, @jsjodin wrote:

In D147172#4240342, @kiranchandramohan wrote:

In D147172#4237283, @jsjodin wrote:

@kiranchandramohan I moved the outlining to the OpenMPIRBuilder, is this more in line of what you were thinking? There is a patch up by @TIFitis for Fortran->Fir lowering https://reviews.llvm.org/D147339. I will prune this patch to only have the OMP+LLVM Dialect -> LLVMIR portions and appropriate tests.

Thanks @jsjodin for the changes. Yes, moving the code to the OpenMPIRBuilder looks fine. I have not looked at the patch in detail. But I see that you have not used the OpenMPIRBuilder's outlining facility (OutlineInfo(s) and the PostOutlineCB). Are there any issues in using these here?

I have reworked the patch a bit to simplify things and added tests. The reason why the OutlineInfos etc are not used for the target regions is probably because the outlining is more complicated. With offloading we have to generate an if-then-else that tries to offload via the openmp runtime and if the runtime fails to launch the kernel the fallback is executed instead.

Also since the outlining is shared with the device codegen, that is also a reason why the OutlineInfos etc is not used.

We have handling for the if clause in both task and parallel. Will the handling in Target be substantially different from these?
Task : D130615
Parallel : D138495 (The latest revision uses a runtime extension, so that we do not have to generate IR for both sequential and parallel version).

The if-then-else is not related to the if clause, it is only for the host fallback. The code in this patch will be updated once the kernel argument handling is implemented for the offloading case. We will have to add the if clause handling as well later on.

mlir/test/Target/LLVMIR/omptarget-llvm.mlir

178

It has to stay because of omp.is_device = #omp.isdevice<is_device = false>

[jsjodin]

Check for unsupported clauses and give error if encountered.

[jdoerfert]

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).

llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h

1828

Nit: newline and docs above; copy&paste error below; it's not the default number but clause value;

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

4166

4180

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.

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

1554

@skatrak You can file a bug. We might be able to handle that fine, but it has not come up yet.

[jsjodin]

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

4180

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.

Should I add the code to build the kernel args, or do you mean empty kernel args?

[jsjodin]

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

4180

I forgot to mention that I am working on a patch to handle the kernel args which adds the kernel launch call. So I could pull that merge that patch into this or keep it separate. @kiranchandramohan, @jdoerfert you have a preference?

[jsjodin]

Fixed naming, braces and comments.

[kiranchandramohan]

Thanks @jsjodin for the changes. Yes, moving the code to the OpenMPIRBuilder looks fine. I have not looked at the patch in detail. But I see that you have not used the OpenMPIRBuilder's outlining facility (OutlineInfo(s) and the PostOutlineCB). Are there any issues in using these here?

I have reworked the patch a bit to simplify things and added tests. The reason why the OutlineInfos etc are not used for the target regions is probably because the outlining is more complicated. With offloading we have to generate an if-then-else that tries to offload via the openmp runtime and if the runtime fails to launch the kernel the fallback is executed instead.

Also since the outlining is shared with the device codegen, that is also a reason why the OutlineInfos etc is not used.

We have handling for the if clause in both task and parallel. Will the handling in Target be substantially different from these?
Task : D130615
Parallel : D138495 (The latest revision uses a runtime extension, so that we do not have to generate IR for both sequential and parallel version).

The if-then-else is not related to the if clause, it is only for the host fallback. The code in this patch will be updated once the kernel argument handling is implemented for the offloading case. We will have to add the if clause handling as well later on.

I was making a general point that, we already have handling in the OpenMPIRBuilder that is capable of conditionally calling a parallel or sequential version. I was hoping that this can be used to generate the if-then-else for the offload.

In general, the concern about not using the OutlineInfo mechanism is whether this will fail to compose with existing. code that uses the OutlineInfo Mechanism.

But if you and @jdoerfert feel this is the right way to go then I don't have any issues.

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

4180

No. I don't have a preference.

[jsjodin]

In D147172#4262744, @kiranchandramohan wrote:

Thanks @jsjodin for the changes. Yes, moving the code to the OpenMPIRBuilder looks fine. I have not looked at the patch in detail. But I see that you have not used the OpenMPIRBuilder's outlining facility (OutlineInfo(s) and the PostOutlineCB). Are there any issues in using these here?

I have reworked the patch a bit to simplify things and added tests. The reason why the OutlineInfos etc are not used for the target regions is probably because the outlining is more complicated. With offloading we have to generate an if-then-else that tries to offload via the openmp runtime and if the runtime fails to launch the kernel the fallback is executed instead.

Also since the outlining is shared with the device codegen, that is also a reason why the OutlineInfos etc is not used.

We have handling for the if clause in both task and parallel. Will the handling in Target be substantially different from these?
Task : D130615
Parallel : D138495 (The latest revision uses a runtime extension, so that we do not have to generate IR for both sequential and parallel version).

The if-then-else is not related to the if clause, it is only for the host fallback. The code in this patch will be updated once the kernel argument handling is implemented for the offloading case. We will have to add the if clause handling as well later on.

I was making a general point that, we already have handling in the OpenMPIRBuilder that is capable of conditionally calling a parallel or sequential version. I was hoping that this can be used to generate the if-then-else for the offload.

In general, the concern about not using the OutlineInfo mechanism is whether this will fail to compose with existing. code that uses the OutlineInfo Mechanism.

But if you and @jdoerfert feel this is the right way to go then I don't have any issues.

I think it can be done.

In D147172#4262744, @kiranchandramohan wrote:

Thanks @jsjodin for the changes. Yes, moving the code to the OpenMPIRBuilder looks fine. I have not looked at the patch in detail. But I see that you have not used the OpenMPIRBuilder's outlining facility (OutlineInfo(s) and the PostOutlineCB). Are there any issues in using these here?

I have reworked the patch a bit to simplify things and added tests. The reason why the OutlineInfos etc are not used for the target regions is probably because the outlining is more complicated. With offloading we have to generate an if-then-else that tries to offload via the openmp runtime and if the runtime fails to launch the kernel the fallback is executed instead.

Also since the outlining is shared with the device codegen, that is also a reason why the OutlineInfos etc is not used.

We have handling for the if clause in both task and parallel. Will the handling in Target be substantially different from these?
Task : D130615
Parallel : D138495 (The latest revision uses a runtime extension, so that we do not have to generate IR for both sequential and parallel version).

The if-then-else is not related to the if clause, it is only for the host fallback. The code in this patch will be updated once the kernel argument handling is implemented for the offloading case. We will have to add the if clause handling as well later on.

I was making a general point that, we already have handling in the OpenMPIRBuilder that is capable of conditionally calling a parallel or sequential version. I was hoping that this can be used to generate the if-then-else for the offload.

In general, the concern about not using the OutlineInfo mechanism is whether this will fail to compose with existing. code that uses the OutlineInfo Mechanism.

But if you and @jdoerfert feel this is the right way to go then I don't have any issues.

I am unsure about how target regions can compose with the other directives that use the OutlineInfo. It might be possible to use the OutlineInfo mechanism, but I don't think it will work with some of the other existing utilities in the OpenMPIRBuilder e.g. emitTargetRegionFunction, which is used in this patch.

[jsjodin]

In D147172#4263033, @jsjodin wrote:

In D147172#4262744, @kiranchandramohan wrote:

Thanks @jsjodin for the changes. Yes, moving the code to the OpenMPIRBuilder looks fine. I have not looked at the patch in detail. But I see that you have not used the OpenMPIRBuilder's outlining facility (OutlineInfo(s) and the PostOutlineCB). Are there any issues in using these here?

I have reworked the patch a bit to simplify things and added tests. The reason why the OutlineInfos etc are not used for the target regions is probably because the outlining is more complicated. With offloading we have to generate an if-then-else that tries to offload via the openmp runtime and if the runtime fails to launch the kernel the fallback is executed instead.

Also since the outlining is shared with the device codegen, that is also a reason why the OutlineInfos etc is not used.

We have handling for the if clause in both task and parallel. Will the handling in Target be substantially different from these?
Task : D130615
Parallel : D138495 (The latest revision uses a runtime extension, so that we do not have to generate IR for both sequential and parallel version).

The if-then-else is not related to the if clause, it is only for the host fallback. The code in this patch will be updated once the kernel argument handling is implemented for the offloading case. We will have to add the if clause handling as well later on.

I was making a general point that, we already have handling in the OpenMPIRBuilder that is capable of conditionally calling a parallel or sequential version. I was hoping that this can be used to generate the if-then-else for the offload.

In general, the concern about not using the OutlineInfo mechanism is whether this will fail to compose with existing. code that uses the OutlineInfo Mechanism.

But if you and @jdoerfert feel this is the right way to go then I don't have any issues.

I think it can be done.

In D147172#4262744, @kiranchandramohan wrote:

Thanks @jsjodin for the changes. Yes, moving the code to the OpenMPIRBuilder looks fine. I have not looked at the patch in detail. But I see that you have not used the OpenMPIRBuilder's outlining facility (OutlineInfo(s) and the PostOutlineCB). Are there any issues in using these here?

I have reworked the patch a bit to simplify things and added tests. The reason why the OutlineInfos etc are not used for the target regions is probably because the outlining is more complicated. With offloading we have to generate an if-then-else that tries to offload via the openmp runtime and if the runtime fails to launch the kernel the fallback is executed instead.

Also since the outlining is shared with the device codegen, that is also a reason why the OutlineInfos etc is not used.

We have handling for the if clause in both task and parallel. Will the handling in Target be substantially different from these?
Task : D130615
Parallel : D138495 (The latest revision uses a runtime extension, so that we do not have to generate IR for both sequential and parallel version).

The if-then-else is not related to the if clause, it is only for the host fallback. The code in this patch will be updated once the kernel argument handling is implemented for the offloading case. We will have to add the if clause handling as well later on.

I was making a general point that, we already have handling in the OpenMPIRBuilder that is capable of conditionally calling a parallel or sequential version. I was hoping that this can be used to generate the if-then-else for the offload.

In general, the concern about not using the OutlineInfo mechanism is whether this will fail to compose with existing. code that uses the OutlineInfo Mechanism.

But if you and @jdoerfert feel this is the right way to go then I don't have any issues.

I am unsure about how target regions can compose with the other directives that use the OutlineInfo. It might be possible to use the OutlineInfo mechanism, but I don't think it will work with some of the other existing utilities in the OpenMPIRBuilder e.g. emitTargetRegionFunction, which is used in this patch.

I think that reusing the offloading utilities is the better option right now. It will allow us to not re-implement a bunch of logic to make faster progress. Also if anything changes with respect to the omp runtime we wouldn't have to modify the code in two places. Since everything will be in the OpenMPIRBuilder we could do some refactoring later of if it seems worth doing.

[kiranchandramohan]

Thanks for the changes. I have a few comments or questions. See inline.

llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h

1835–1837

Nit: Please add a full stop/dot at the end for all the three above to keep it consistent.

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

4118

Nit: Is 4 significant here?

4119–4121

Nit: braces not required.

4120

In the past we have seen that if the number of live-ins/parameters go above a certain number that leads to ABI issues on some platforms.
The OpenMPIRBuilder flow packs the arguments into a struct and then unpacks it in the outlined function. Is this not a concern for offloading?

4141–4142

Nit: spell the types.

4144–4151

Nit: braces not required.

4196

Nit: Is this leftover debug code?

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

1571

Is this required? If so, please add it to other if statements as well.

1602

Return success or failure?

1612–1614

Nit: braces not required.

mlir/test/Target/LLVMIR/omptarget-llvm.mlir

177

Please also add a trivial test for an OpenMP construct that can occur inside the target region, just to ensure that the target region can contain other constructs and it works OK.

Something very simple like the following.

omp.target {
  omp.parallel {
     ....
     omp.terminator
  }
  omp.terminator
}

178

Nit: If the dlti is not required for this patch, then you can consider omitting it.

[jsjodin]

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

4118

Nit: Is 4 significant here?

No, I picked something small that hopefully doesn't require a resize. Not sure if there is a better value to pick?

4120

In the past we have seen that if the number of live-ins/parameters go above a certain number that leads to ABI issues on some platforms.
The OpenMPIRBuilder flow packs the arguments into a struct and then unpacks it in the outlined function. Is this not a concern for offloading?

Clang does not pack arguments into a struct for the outlined fallback (host) for target regions. It could be that the number of parameters are fewer compared to other outlining cases.
It might be desirable to use the kernel args structure so both the host and device use the same argument passing, but this is not done today.

4196

Yes, thanks for finding it!

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

1571

Yes, I think it is okay to error out. I will add it to the other cases.

1602

Success for now. Basically ignore this so we can get through the driver. There is a follow up patch https://reviews.llvm.org/D147940 that will allow us to enable device codegen.

mlir/test/Target/LLVMIR/omptarget-llvm.mlir

177

Please also add a trivial test for an OpenMP construct that can occur inside the target region, just to ensure that the target region can contain other constructs and it works OK.

Something very simple like the following.

omp.target {
  omp.parallel {
     ....
     omp.terminator
  }
  omp.terminator
}

Sure, I can add a test for this.

178

Nit: If the dlti is not required for this patch, then you can consider omitting it.

We can remove dlti it for this patch.

[kiranchandramohan]

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

4118

Doesn't the programmers manual recommend omitting if the value is not significant?

In the absence of a well-motivated choice for the number of inlined elements N, it is recommended to use SmallVector<T> (that is, omitting the N). This will choose a default number of inlined elements reasonable for allocation on the stack (for example, trying to keep sizeof(SmallVector<T>) around 64 bytes).

https://llvm.org/docs/ProgrammersManual.html#llvm-adt-smallvector-h

4120

OK. It is fine for now.

[jsjodin]

Fixed comments. Added test with omp.parallel.

[kiranchandramohan]

A few comments/questions.

mlir/lib/Target/LLVMIR/CMakeLists.txt

60

Is this required?

87

Is this required?

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

1643–1644

Can this function be sunk into the OpenMPIRBuilder? If it is just a call like this then please consider, otherwise ignore.

mlir/lib/Target/LLVMIR/ModuleTranslation.cpp

1348–1356

Ideally, this code should be in OpenMPToLLVMIRTranslation.cpp. Can this be moved there?

[jsjodin]

mlir/lib/Target/LLVMIR/CMakeLists.txt

60

Is this required?

No, thanks for pointing this out. I think it was needed in an earlier iteration.

87

Is this required?

Not needed here either. Thanks!

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

1643–1644

Can this function be sunk into the OpenMPIRBuilder? If it is just a call like this then please consider, otherwise ignore.

No, unfortunately it cannot be replaced with a simple function call.

mlir/lib/Target/LLVMIR/ModuleTranslation.cpp

1348–1356

Ideally, this code should be in OpenMPToLLVMIRTranslation.cpp. Can this be moved there?

I'm not sure how to do that. The code in OpenMPToLLVMIRTranslation code implements an interface and does not have an internal state, and we probably don't want to set the config in every function. The object that holds the state is ModuleTranslation. We could move this code to the constructor of the ModuleTranslation class if you prefer.

[jsjodin]

Removed useless dependencies in cmake file. Moved setting of config to the ModuleTranslation constructor.

[kiranchandramohan]

mlir/lib/Target/LLVMIR/ModuleTranslation.cpp

460

This will always create the OpenMPIRBuilder, even for non-OpenMP flows which is not what we want.

getOpenMPBuilder creates the Builder if it is not yet created. We can consider setting the config there.

[jsjodin]

Moved setting of the OMPIRBuilder cofiguration to getOpenMPBuilder.

[jsjodin]

mlir/lib/Target/LLVMIR/ModuleTranslation.cpp

460

This will always create the OpenMPIRBuilder, even for non-OpenMP flows which is not what we want.

getOpenMPBuilder creates the Builder if it is not yet created. We can consider setting the config there.

Yes, that is a much better option. Moved the code there instead.

[jsjodin]

mlir/lib/Target/LLVMIR/ModuleTranslation.cpp

1270

I will remove this recursive call.

[jsjodin]

Removed recursive call + merge.

[kiranchandramohan]

LGTM. Please check with @jdoerfert before you submit.

It will be good if someone in your team can go through the tests in detail.

mlir/lib/Target/LLVMIR/CMakeLists.txt

42

Nit: Please reconfirm that this is required.

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

14

Is this include required?

[jsjodin]

mlir/lib/Target/LLVMIR/CMakeLists.txt

42

It is needed, link error if not there.

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

14

Not needed, will remove.

[jsjodin]

Remove useless include

[jsjodin]

In D147172#4295771, @kiranchandramohan wrote:

LGTM. Please check with @jdoerfert before you submit.

It will be good if someone in your team can go through the tests in detail.

Thank you for taking the time to review this patch!
I will let @jdoerfert review, and have someone look at the test.

[jdoerfert]

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

4118

[jsjodin]

Remove 'llvm::'

[agozillon]

Looked over everything one final time, including the tests, it all LGTM.

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

1581

I have a patch that moves this into the OMP IR builder: https://reviews.llvm.org/D149162 so it might be possible to replace this at some point in the near future with it! It's used for declare target variables as well, and perhaps some other things. I can likely create a patch to change it to utilise the OMPIRBuilder version if/when it lands. I bring it up as a future TODO, for myself (or whomever) rather than something to do in this patch.

[kstoimenov]

Looks like this broke our sanitizer bot: https://lab.llvm.org/buildbot/#/builders/5/builds/33303. Could you please revert or fix forward?

Thanks!

[Mogball]

mlir/lib/Target/LLVMIR/CMakeLists.txt

42

This seems problematic. OpenMP dialect should not be a dependency on the generic LLVMIR export. It seems there is a layering issue here. You need to add your own LLVMIR export target with the OpenMP dialect registered.

[kiranchandramohan]

mlir/lib/Target/LLVMIR/CMakeLists.txt

42

@jsjodin If the link error is coming due to the use of the OffloadInterface then you can remove its usage and access attributes directly. If accessing attributes still causes a link error then you can consider initializing the OpenMPIRBuilder using the amdendOperation flow. You might have to move convertModule to somewhere earlier for this to work. If that is also not possible and there are no other ideas then please revert the patch.

[jsjodin]

mlir/lib/Target/LLVMIR/CMakeLists.txt

42

Sorry for the delayed response. I am looking into this issue.

[jsjodin]

mlir/lib/Target/LLVMIR/CMakeLists.txt

42

A fix is in this patch: https://reviews.llvm.org/D151745

[jsjodin]

mlir/lib/Target/LLVMIR/CMakeLists.txt

42

A fix is in this patch: https://reviews.llvm.org/D151745

This was committed now.

llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h

public:
   OpenMPIRBuilder::InsertPointTy createTargetData(
       const LocationDescription &Loc, OpenMPIRBuilder::InsertPointTy CodeGenIP,
       SmallVectorImpl<uint64_t> &MapTypeFlags,
       SmallVectorImpl<Constant *> &MapNames,
       struct MapperAllocas &MapperAllocas, bool IsBegin, int64_t DeviceID,
       Value *IfCond, BodyGenCallbackTy ProcessMapOpCB,
       BodyGenCallbackTy BodyGenCB = {});
 
+  using TargetBodyGenCallbackTy = function_ref<InsertPointTy(
+      InsertPointTy AllocaIP, InsertPointTy CodeGenIP)>;
+
+  /// Generator for '#omp target'
+  ///
+  /// \param Loc where the target data construct was encountered.
+  /// \param CodeGenIP The insertion point where the call to the outlined
+  /// function should be emitted.
+  /// \param EntryInfo The entry information about the function.
+  /// \param NumTeams Number of teams specified in the num_teams clause.
+  /// \param NumThreads Number of teams specified in the thread_limit clause.
+  /// \param Inputs The input values to the region that will be passed.
+  /// as arguments to the outlined function.
+  /// \param BodyGenCB Callback that will generate the region code.
+  InsertPointTy createTarget(const LocationDescription &Loc,
+                             OpenMPIRBuilder::InsertPointTy CodeGenIP,
+                             TargetRegionEntryInfo &EntryInfo, int32_t NumTeams,
+                             int32_t NumThreads,
+                             SmallVectorImpl<Value *> &Inputs,
+                             TargetBodyGenCallbackTy BodyGenCB);
+
   /// Declarations for LLVM-IR types (simple, array, function and structure) are
   /// generated below. Their names are defined and used in OpenMPKinds.def. Here
   /// we provide the declarations, the initializeTypes function will provide the
   /// values.
   ///
   ///{
 #define OMP_TYPE(VarName, InitValue) Type *VarName = nullptr;
 #define OMP_ARRAY_TYPE(VarName, ElemTy, ArraySize)                             \

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

OpenMPIRBuilder::InsertPointTy OpenMPIRBuilder::createTargetData(
   // Update the insertion point and remove the terminator we introduced.
   Builder.SetInsertPoint(UI->getParent());
   if (IfCond)
     UI->getParent()->setName("omp_if.end");
   UI->eraseFromParent();
   return Builder.saveIP();
 }
 
+static Function *
+createOutlinedFunction(IRBuilderBase &Builder, StringRef FuncName,
+                       SmallVectorImpl<Value *> &Inputs,
+                       OpenMPIRBuilder::TargetBodyGenCallbackTy &CBFunc) {
+  llvm::SmallVector<llvm::Type *> ParameterTypes;
+  for (auto &Arg : Inputs)
+    ParameterTypes.push_back(Arg->getType());
+
+  auto FuncType = FunctionType::get(Builder.getVoidTy(), ParameterTypes,
+                                    /*isVarArg*/ false);
+  auto Func = Function::Create(FuncType, GlobalValue::InternalLinkage, FuncName,
+                               Builder.GetInsertBlock()->getModule());
+
+  // Save insert point.
+  auto OldInsertPoint = Builder.saveIP();
+
+  // Generate the region into the function.
+  BasicBlock *EntryBB = BasicBlock::Create(Builder.getContext(), "entry", Func);
+  Builder.SetInsertPoint(EntryBB);
+  Builder.restoreIP(CBFunc(Builder.saveIP(), Builder.saveIP()));
+
+  // Insert return instruction.
+  Builder.CreateRetVoid();
+
+  // Rewrite uses of input valus to parameters.
+  for (auto InArg : zip(Inputs, Func->args())) {
+    Value *Input = std::get<0>(InArg);
+    Argument &Arg = std::get<1>(InArg);
+
+    // Collect all the instructions
+    for (User *User : make_early_inc_range(Input->users()))
+      if (auto Instr = dyn_cast<Instruction>(User))
+        if (Instr->getFunction() == Func)
+          Instr->replaceUsesOfWith(Input, &Arg);
+  }
+
+  // Restore insert point.
+  Builder.restoreIP(OldInsertPoint);
+
+  return Func;
+}
+
+static void
+emitTargetOutlinedFunction(OpenMPIRBuilder &OMPBuilder, IRBuilderBase &Builder,
+                           TargetRegionEntryInfo &EntryInfo,
+                           Function *&OutlinedFn, int32_t NumTeams,
+                           int32_t NumThreads, SmallVectorImpl<Value *> &Inputs,
+                           OpenMPIRBuilder::TargetBodyGenCallbackTy &CBFunc) {
+
+  OpenMPIRBuilder::FunctionGenCallback &&GenerateOutlinedFunction =
+      [&Builder, &Inputs, &CBFunc](StringRef EntryFnName) {
+        return createOutlinedFunction(Builder, EntryFnName, Inputs, CBFunc);
+      };
+
+  Constant *OutlinedFnID;
+  OMPBuilder.emitTargetRegionFunction(EntryInfo, GenerateOutlinedFunction,
+                                      NumTeams, NumThreads, true, OutlinedFn,
+                                      OutlinedFnID);
+}
+
+static void emitTargetCall(IRBuilderBase &Builder, Function *OutlinedFn,
+                           SmallVectorImpl<Value *> &Args) {
+  // TODO: Add kernel launch call when device codegen is supported.
+  Builder.CreateCall(OutlinedFn, Args);
+}
+
+OpenMPIRBuilder::InsertPointTy OpenMPIRBuilder::createTarget(
+    const LocationDescription &Loc, OpenMPIRBuilder::InsertPointTy CodeGenIP,
+    TargetRegionEntryInfo &EntryInfo, int32_t NumTeams, int32_t NumThreads,
+    SmallVectorImpl<Value *> &Args, TargetBodyGenCallbackTy CBFunc) {
+  if (!updateToLocation(Loc))
+    return InsertPointTy();
+
+  Builder.restoreIP(CodeGenIP);
+
+  Function *OutlinedFn;
+  emitTargetOutlinedFunction(*this, Builder, EntryInfo, OutlinedFn, NumTeams,
+                             NumThreads, Args, CBFunc);
+  emitTargetCall(Builder, OutlinedFn, Args);
+  return Builder.saveIP();
+}
+
 std::string OpenMPIRBuilder::getNameWithSeparators(ArrayRef<StringRef> Parts,
                                                    StringRef FirstSeparator,
                                                    StringRef Separator) {
   SmallString<128> Buffer;
   llvm::raw_svector_ostream OS(Buffer);
   StringRef Sep = FirstSeparator;
   for (StringRef Part : Parts) {
     OS << Sep << Part;
     Sep = Separator;

llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp

TEST_F(OpenMPIRBuilderTest, TargetDataRegion) {
   EXPECT_TRUE(TargetDataCall->getOperand(1)->getType()->isIntegerTy(64));
   EXPECT_TRUE(TargetDataCall->getOperand(2)->getType()->isIntegerTy(32));
   EXPECT_TRUE(TargetDataCall->getOperand(8)->getType()->isPointerTy());
 
   Builder.CreateRetVoid();
   EXPECT_FALSE(verifyModule(*M, &errs()));
 }
 
+TEST_F(OpenMPIRBuilderTest, TargetRegion) {
+  using InsertPointTy = OpenMPIRBuilder::InsertPointTy;
+  OpenMPIRBuilder OMPBuilder(*M);
+  OMPBuilder.initialize();
+  OpenMPIRBuilderConfig Config(false, false, false, false);
+  OMPBuilder.setConfig(Config);
+  F->setName("func");
+  IRBuilder<> Builder(BB);
+  auto Int32Ty = Builder.getInt32Ty();
+
+  AllocaInst *APtr = Builder.CreateAlloca(Int32Ty, nullptr, "a_ptr");
+  AllocaInst *BPtr = Builder.CreateAlloca(Int32Ty, nullptr, "b_ptr");
+  AllocaInst *CPtr = Builder.CreateAlloca(Int32Ty, nullptr, "c_ptr");
+
+  Builder.CreateStore(Builder.getInt32(10), APtr);
+  Builder.CreateStore(Builder.getInt32(20), BPtr);
+  auto BodyGenCB = [&](InsertPointTy AllocaIP,
+                       InsertPointTy CodeGenIP) -> InsertPointTy {
+    Builder.restoreIP(CodeGenIP);
+    LoadInst *AVal = Builder.CreateLoad(Int32Ty, APtr);
+    LoadInst *BVal = Builder.CreateLoad(Int32Ty, BPtr);
+    Value *Sum = Builder.CreateAdd(AVal, BVal);
+    Builder.CreateStore(Sum, CPtr);
+    return Builder.saveIP();
+  };
+
+  llvm::SmallVector<llvm::Value *> Inputs;
+  Inputs.push_back(APtr);
+  Inputs.push_back(BPtr);
+  Inputs.push_back(CPtr);
+
+  TargetRegionEntryInfo EntryInfo("func", 42, 4711, 17);
+  OpenMPIRBuilder::LocationDescription OmpLoc({Builder.saveIP(), DL});
+  Builder.restoreIP(OMPBuilder.createTarget(OmpLoc, Builder.saveIP(), EntryInfo,
+                                            -1, -1, Inputs, BodyGenCB));
+  OMPBuilder.finalize();
+  Builder.CreateRetVoid();
+
+  // Check the outlined call
+  auto Iter = F->getEntryBlock().rbegin();
+  CallInst *Call = dyn_cast<CallInst>(&*(++Iter));
+  EXPECT_NE(Call, nullptr);
+
+  // Check that the correct aguments are passed in
+  for (auto ArgInput : zip(Call->args(), Inputs)) {
+    EXPECT_EQ(std::get<0>(ArgInput), std::get<1>(ArgInput));
+  }
+
+  // Check that the outlined function exists with the expected prefix
+  Function *OutlinedFunc = Call->getCalledFunction();
+  EXPECT_NE(OutlinedFunc, nullptr);
+  StringRef FunctionName = OutlinedFunc->getName();
+  EXPECT_TRUE(FunctionName.startswith("__omp_offloading"));
+  EXPECT_FALSE(verifyModule(*M, &errs()));
+}
+
 TEST_F(OpenMPIRBuilderTest, CreateTask) {
   using InsertPointTy = OpenMPIRBuilder::InsertPointTy;
   OpenMPIRBuilder OMPBuilder(*M);
   OMPBuilder.initialize();
   F->setName("func");
   IRBuilder<> Builder(BB);
 
   AllocaInst *ValPtr32 = Builder.CreateAlloca(Builder.getInt32Ty());

mlir/include/mlir/Target/LLVMIR/ModuleTranslation.h

public:
   /// Finds an LLVM IR global value that corresponds to the given MLIR operation
   /// defining a global value.
   llvm::GlobalValue *lookupGlobal(Operation *op) {
     return globalsMapping.lookup(op);
   }
 
   /// Returns the OpenMP IR builder associated with the LLVM IR module being
   /// constructed.
-  llvm::OpenMPIRBuilder *getOpenMPBuilder() {
-    if (!ompBuilder) {
-      ompBuilder = std::make_unique<llvm::OpenMPIRBuilder>(*llvmModule);
-      ompBuilder->initialize();
-    }
-    return ompBuilder.get();
-  }
+  llvm::OpenMPIRBuilder *getOpenMPBuilder();
 
   /// Translates the given location.
   const llvm::DILocation *translateLoc(Location loc, llvm::DILocalScope *scope);
 
   /// Translates the given LLVM debug info metadata.
   llvm::Metadata *translateDebugInfo(LLVM::DINodeAttr attr);
 
   /// Translates the contents of the given block to LLVM IR using this

mlir/lib/Target/LLVMIR/CMakeLists.txt

add_mlir_translation_library(MLIRTargetLLVMIRExport
   TransformUtils
   TargetParser
 
   LINK_LIBS PUBLIC
   MLIRDLTIDialect
   MLIRLLVMDialect
   MLIRLLVMIRTransforms
   MLIRTranslateLib
+  MLIROpenMPDialect
   )
 
 add_mlir_translation_library(MLIRToLLVMIRTranslationRegistration
   ConvertToLLVMIR.cpp
 
   LINK_LIBS PUBLIC
   MLIRArmNeonToLLVMIRTranslation
   MLIRArmSVEToLLVMIRTranslation
   MLIRAMXToLLVMIRTranslation
   MLIRBuiltinToLLVMIRTranslation
   MLIRGPUToLLVMIRTranslation
   MLIRX86VectorToLLVMIRTranslation
   MLIRLLVMToLLVMIRTranslation
   MLIRNVVMToLLVMIRTranslation
   MLIROpenACCToLLVMIRTranslation
   MLIROpenMPToLLVMIRTranslation
   MLIRROCDLToLLVMIRTranslation
   )
 
 add_mlir_translation_library(MLIRTargetLLVMIRImport
   DataLayoutImporter.cpp
   DebugImporter.cpp
   LoopAnnotationImporter.cpp
   ModuleImport.cpp
   TypeFromLLVM.cpp
 
add_mlir_translation_library(MLIRTargetLLVMIRImport
   MLIRTranslateLib
   )
 
 add_mlir_translation_library(MLIRFromLLVMIRTranslationRegistration
   ConvertFromLLVMIR.cpp
 
   LINK_LIBS PUBLIC
   MLIRLLVMIRToLLVMTranslation
   )

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

 //===- OpenMPToLLVMIRTranslation.cpp - Translate OpenMP dialect to LLVM IR-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements a translation between the MLIR OpenMP dialect and LLVM
 // IR.
 //
 //===----------------------------------------------------------------------===//
 #include "mlir/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.h"
+#include "mlir/Dialect/LLVMIR/LLVMOps.cpp.inc"
 #include "mlir/Dialect/OpenMP/OpenMPDialect.h"
+#include "mlir/Dialect/OpenMP/OpenMPInterfaces.h"
 #include "mlir/IR/IRMapping.h"
 #include "mlir/IR/Operation.h"
 #include "mlir/Support/LLVM.h"
 #include "mlir/Target/LLVMIR/Dialect/OpenMPCommon.h"
 #include "mlir/Target/LLVMIR/ModuleTranslation.h"
+#include "mlir/Transforms/RegionUtils.h"
 
 #include "llvm/ADT/SetVector.h"
 #include "llvm/ADT/TypeSwitch.h"
 #include "llvm/Frontend/OpenMP/OMPConstants.h"
 #include "llvm/Frontend/OpenMP/OMPIRBuilder.h"
 #include "llvm/IR/DebugInfoMetadata.h"
 #include "llvm/IR/IRBuilder.h"
+#include "llvm/Support/FileSystem.h"
 
 using namespace mlir;
 
 namespace {
 static llvm::omp::ScheduleKind
 convertToScheduleKind(std::optional<omp::ClauseScheduleKind> schedKind) {
   if (!schedKind.has_value())
     return llvm::omp::OMP_SCHEDULE_Default;
convertOmpTargetData(Operation *op, llvm::IRBuilderBase &builder,
   return bodyGenStatus;
 }
 
 /// Lowers the FlagsAttr which is applied to the module on the device
 /// pass when offloading, this attribute contains OpenMP RTL globals that can
 /// be passed as flags to the frontend, otherwise they are set to default
 LogicalResult convertFlagsAttr(Operation *op, mlir::omp::FlagsAttr attribute,
                                LLVM::ModuleTranslation &moduleTranslation) {
   llvm::OpenMPIRBuilder *ompBuilder = moduleTranslation.getOpenMPBuilder();
 
   ompBuilder->createGlobalFlag(
       attribute.getDebugKind() /*LangOpts().OpenMPTargetDebug*/,
       "__omp_rtl_debug_kind");
   ompBuilder->createGlobalFlag(
       attribute
           .getAssumeTeamsOversubscription() /*LangOpts().OpenMPTeamSubscription*/
       ,
       "__omp_rtl_assume_teams_oversubscription");
   ompBuilder->createGlobalFlag(
       attribute
           .getAssumeThreadsOversubscription() /*LangOpts().OpenMPThreadSubscription*/
       ,
       "__omp_rtl_assume_threads_oversubscription");
   ompBuilder->createGlobalFlag(
       attribute.getAssumeNoThreadState() /*LangOpts().OpenMPNoThreadState*/,
       "__omp_rtl_assume_no_thread_state");
   ompBuilder->createGlobalFlag(
       attribute
           .getAssumeNoNestedParallelism() /*LangOpts().OpenMPNoNestedParallelism*/
       ,
       "__omp_rtl_assume_no_nested_parallelism");
 
   return success();
 }
 
+static llvm::TargetRegionEntryInfo
+getTargetEntryUniqueInfo(omp::TargetOp targetOp,
+                         llvm::StringRef parentName = "") {
+  auto fileLoc = targetOp.getLoc()->findInstanceOf<FileLineColLoc>();
+
+  assert(fileLoc && "No file found from location");
+  StringRef fileName = fileLoc.getFilename().getValue();
+
+  llvm::sys::fs::UniqueID id;
+  if (auto ec = llvm::sys::fs::getUniqueID(fileName, id)) {
+    targetOp.emitError("Unable to get unique ID for file");
+  }
+
+  uint64_t line = fileLoc.getLine();
+  return llvm::TargetRegionEntryInfo(parentName, id.getDevice(), id.getFile(),
+                                     line);
+}
+
+static bool targetOpSupported(Operation &opInst) {
+  auto targetOp = cast<omp::TargetOp>(opInst);
+  if (targetOp.getIfExpr()) {
+    opInst.emitError("If clause not yet supported");
+    return false;
+  }
+
+  if (targetOp.getDevice()) {
+    opInst.emitError("Device clause not yet supported");
+    return false;
+  }
+
+  if (targetOp.getThreadLimit()) {
+    opInst.emitError("Thread limit clause not yet supported");
+    return false;
+  }
+
+  if (targetOp.getNowait()) {
+    opInst.emitError("Nowait clause not yet supported");
+    return false;
+  }
+
+  return true;
+}
+
+static LogicalResult
+convertOmpTarget(Operation &opInst, llvm::IRBuilderBase &builder,
+                 LLVM::ModuleTranslation &moduleTranslation) {
+
+  if (!targetOpSupported(opInst))
+    return failure();
+
+  bool isDevice = false;
+  if (auto offloadMod = dyn_cast<mlir::omp::OffloadModuleInterface>(
+          opInst.getParentOfType<mlir::ModuleOp>().getOperation())) {
+    isDevice = offloadMod.getIsDevice();
+  }
+
+  if (isDevice) // TODO: Implement device codegen.
+    return success();
+
+  auto targetOp = cast<omp::TargetOp>(opInst);
+  auto &targetRegion = targetOp.getRegion();
+
+  llvm::SetVector<Value> operandSet;
+  getUsedValuesDefinedAbove(targetRegion, operandSet);
+
+  // Collect the input arguments.
+  llvm::SmallVector<llvm::Value *> inputs;
+  for (Value operand : operandSet)
+    inputs.push_back(moduleTranslation.lookupValue(operand));
+
+  LogicalResult bodyGenStatus = success();
+
+  using InsertPointTy = llvm::OpenMPIRBuilder::InsertPointTy;
+  auto bodyCB = [&](InsertPointTy allocaIP,
+                    InsertPointTy codeGenIP) -> InsertPointTy {
+    builder.restoreIP(codeGenIP);
+    llvm::BasicBlock *exitBlock = convertOmpOpRegions(
+        targetRegion, "omp.target", builder, moduleTranslation, bodyGenStatus);
+    builder.SetInsertPoint(exitBlock);
+    return builder.saveIP();
+  };
+
+  llvm::OpenMPIRBuilder::LocationDescription ompLoc(builder);
+  StringRef parentName = opInst.getParentOfType<LLVM::LLVMFuncOp>().getName();
+  llvm::TargetRegionEntryInfo entryInfo =
+      getTargetEntryUniqueInfo(targetOp, parentName);
+  int32_t defaultValTeams = -1;
+  int32_t defaultValThreads = -1;
+
+  builder.restoreIP(moduleTranslation.getOpenMPBuilder()->createTarget(
+      ompLoc, builder.saveIP(), entryInfo, defaultValTeams, defaultValThreads,
+      inputs, bodyCB));
+
+  return bodyGenStatus;
+}
+
 namespace {
 
 /// Implementation of the dialect interface that converts operations belonging
 /// to the OpenMP dialect to LLVM IR.
 class OpenMPDialectLLVMIRTranslationInterface
     : public LLVMTranslationDialectInterface {
 public:
   using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
return llvm::TypeSwitch<Operation *, LogicalResult>(op)
         return success();
       })
       .Case([&](omp::ThreadprivateOp) {
         return convertOmpThreadprivate(*op, builder, moduleTranslation);
       })
       .Case<omp::DataOp, omp::EnterDataOp, omp::ExitDataOp>([&](auto op) {
         return convertOmpTargetData(op, builder, moduleTranslation);
       })
+      .Case([&](omp::TargetOp) {
+        return convertOmpTarget(*op, builder, moduleTranslation);
+      })
       .Default([&](Operation *inst) {
         return inst->emitError("unsupported OpenMP operation: ")
                << inst->getName();
       });
 }
 
 void mlir::registerOpenMPDialectTranslation(DialectRegistry &registry) {
   registry.insert<omp::OpenMPDialect>();

mlir/lib/Target/LLVMIR/ModuleTranslation.cpp

 #include "AttrKindDetail.h"
 #include "DebugTranslation.h"
 #include "LoopAnnotationTranslation.h"
 #include "mlir/Dialect/DLTI/DLTI.h"
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/Dialect/LLVMIR/LLVMInterfaces.h"
 #include "mlir/Dialect/LLVMIR/Transforms/LegalizeForExport.h"
 #include "mlir/Dialect/OpenMP/OpenMPDialect.h"
+#include "mlir/Dialect/OpenMP/OpenMPInterfaces.h"
 #include "mlir/IR/Attributes.h"
 #include "mlir/IR/BuiltinOps.h"
 #include "mlir/IR/BuiltinTypes.h"
 #include "mlir/IR/RegionGraphTraits.h"
 #include "mlir/Support/LLVM.h"
 #include "mlir/Target/LLVMIR/LLVMTranslationInterface.h"
 #include "mlir/Target/LLVMIR/TypeToLLVM.h"
 
: mlirModule(module), llvmModule(std::move(llvmModule)),
           std::make_unique<DebugTranslation>(module, *this->llvmModule)),
       loopAnnotationTranslation(std::make_unique<LoopAnnotationTranslation>(
           module, *this->llvmModule)),
       typeTranslator(this->llvmModule->getContext()),
       iface(module->getContext()) {
   assert(satisfiesLLVMModule(mlirModule) &&
          "mlirModule should honor LLVM's module semantics.");
 }
+
 ModuleTranslation::~ModuleTranslation() {
   if (ompBuilder)
     ompBuilder->finalize();
 }
 
 void ModuleTranslation::forgetMapping(Region &region) {
   SmallVector<Region *> toProcess;
   toProcess.push_back(&region);
   while (!toProcess.empty()) {
     Region *current = toProcess.pop_back_val();
     for (Block &block : *current) {
       blockMapping.erase(&block);
       for (Value arg : block.getArguments())
         valueMapping.erase(arg);
       for (Operation &op : block) {

 SmallVector<llvm::Value *> ModuleTranslation::lookupValues(ValueRange values) {
   SmallVector<llvm::Value *> remapped;
   remapped.reserve(values.size());
   for (Value v : values)
     remapped.push_back(lookupValue(v));
   return remapped;
 }
 
+llvm::OpenMPIRBuilder *ModuleTranslation::getOpenMPBuilder() {
+  if (!ompBuilder) {
+    ompBuilder = std::make_unique<llvm::OpenMPIRBuilder>(*llvmModule);
+    ompBuilder->initialize();
+
+    bool isDevice = false;
+    if (auto offloadMod =
+            dyn_cast<mlir::omp::OffloadModuleInterface>(mlirModule))
+      isDevice = offloadMod.getIsDevice();
+
+    // TODO: set the flags when available
+    llvm::OpenMPIRBuilderConfig Config(
+        isDevice, /* IsTargetCodegen */ false,
+        /* HasRequiresUnifiedSharedMemory */ false,
+        /* OpenMPOffloadMandatory */ false);
+    getOpenMPBuilder()->setConfig(Config);
+  }
+  return ompBuilder.get();
+}
+
 const llvm::DILocation *
 ModuleTranslation::translateLoc(Location loc, llvm::DILocalScope *scope) {
   return debugTranslation->translateLoc(loc, scope);
 }
 
 llvm::Metadata *ModuleTranslation::translateDebugInfo(LLVM::DINodeAttr attr) {
   return debugTranslation->translate(attr);
 }
std::unique_ptr<llvm::Module> llvmModule =
       prepareLLVMModule(module, llvmContext, name);
   if (!llvmModule)
     return nullptr;
 
   LLVM::ensureDistinctSuccessors(module);
 
   ModuleTranslation translator(module, std::move(llvmModule));
   if (failed(translator.convertFunctionSignatures()))
     return nullptr;
   if (failed(translator.convertGlobals()))
     return nullptr;
   if (failed(translator.createAccessGroupMetadata()))
     return nullptr;
   if (failed(translator.createAliasScopeMetadata()))
     return nullptr;
   if (failed(translator.createTBAAMetadata()))
     return nullptr;
   if (failed(translator.convertFunctions()))
     return nullptr;
 
   // Convert other top-level operations if possible.
   llvm::IRBuilder<> llvmBuilder(llvmContext);
   for (Operation &o : getModuleBody(module).getOperations()) {
     if (!isa<LLVM::LLVMFuncOp, LLVM::GlobalOp, LLVM::GlobalCtorsOp,
              LLVM::GlobalDtorsOp, LLVM::MetadataOp>(&o) &&

mlir/test/Target/LLVMIR/omptarget-llvm.mlir

 // CHECK:         %[[VAL_38:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_1]], i32 0, i32 0
 // CHECK:         %[[VAL_39:.*]] = getelementptr inbounds [2 x i64], ptr %[[VAL_2]], i32 0, i32 0
 // CHECK:         call void @__tgt_target_data_end_mapper(ptr @3, i64 -1, i32 2, ptr %[[VAL_37]], ptr %[[VAL_38]], ptr %[[VAL_39]], ptr @.offload_maptypes.1, ptr @.offload_mapnames.2, ptr null)
 // CHECK:         br label %[[VAL_30]]
 // CHECK:       omp_if.end5:                                      ; preds = %[[VAL_14]], %[[VAL_29]]
 // CHECK:         ret void
 
 // -----
+
+module attributes {omp.is_device = #omp.isdevice<is_device = false>} {
+  llvm.func @omp_target_region_() {
+    %0 = llvm.mlir.constant(20 : i32) : i32
+    %1 = llvm.mlir.constant(10 : i32) : i32
+    %2 = llvm.mlir.constant(1 : i64) : i64
+    %3 = llvm.alloca %2 x i32 {bindc_name = "a", in_type = i32, operand_segment_sizes = array<i32: 0, 0>, uniq_name = "_QFomp_target_regionEa"} : (i64) -> !llvm.ptr<i32>
+    %4 = llvm.mlir.constant(1 : i64) : i64
+    %5 = llvm.alloca %4 x i32 {bindc_name = "b", in_type = i32, operand_segment_sizes = array<i32: 0, 0>, uniq_name = "_QFomp_target_regionEb"} : (i64) -> !llvm.ptr<i32>
+    %6 = llvm.mlir.constant(1 : i64) : i64
+    %7 = llvm.alloca %6 x i32 {bindc_name = "c", in_type = i32, operand_segment_sizes = array<i32: 0, 0>, uniq_name = "_QFomp_target_regionEc"} : (i64) -> !llvm.ptr<i32>
+    llvm.store %1, %3 : !llvm.ptr<i32>
+    llvm.store %0, %5 : !llvm.ptr<i32>
+    omp.target   {
+      %8 = llvm.load %3 : !llvm.ptr<i32>
+      %9 = llvm.load %5 : !llvm.ptr<i32>
+      %10 = llvm.add %8, %9  : i32
+      llvm.store %10, %7 : !llvm.ptr<i32>
+      omp.terminator
+    }
+    llvm.return
+  }
+}
+
+// CHECK: call void @__omp_offloading_[[DEV:.*]]_[[FIL:.*]]_omp_target_region__l[[LINE:.*]](ptr %{{.*}}, ptr %{{.*}}, ptr %{{.*}})
+// CHECK: define internal void @__omp_offloading_[[DEV]]_[[FIL]]_omp_target_region__l[[LINE]](ptr %[[ADDR_A:.*]], ptr %[[ADDR_B:.*]], ptr %[[ADDR_C:.*]])
+// CHECK: %[[VAL_A:.*]] = load i32, ptr %[[ADDR_A]], align 4
+// CHECK: %[[VAL_B:.*]] = load i32, ptr %[[ADDR_B]], align 4
+// CHECK: %[[SUM:.*]] = add i32 %[[VAL_A]], %[[VAL_B]]
+// CHECK: store i32 %[[SUM]], ptr %[[ADDR_C]], align 4
+
+// -----
+
+module attributes {omp.is_device = #omp.isdevice<is_device = false>} {
+  llvm.func @omp_target_region_() {
+    %0 = llvm.mlir.constant(20 : i32) : i32
+    %1 = llvm.mlir.constant(10 : i32) : i32
+    %2 = llvm.mlir.constant(1 : i64) : i64
+    %3 = llvm.alloca %2 x i32 {bindc_name = "a", in_type = i32, operand_segment_sizes = array<i32: 0, 0>, uniq_name = "_QFomp_target_regionEa"} : (i64) -> !llvm.ptr<i32>
+    %4 = llvm.mlir.constant(1 : i64) : i64
+    %5 = llvm.alloca %4 x i32 {bindc_name = "b", in_type = i32, operand_segment_sizes = array<i32: 0, 0>, uniq_name = "_QFomp_target_regionEb"} : (i64) -> !llvm.ptr<i32>
+    %6 = llvm.mlir.constant(1 : i64) : i64
+    %7 = llvm.alloca %6 x i32 {bindc_name = "c", in_type = i32, operand_segment_sizes = array<i32: 0, 0>, uniq_name = "_QFomp_target_regionEc"} : (i64) -> !llvm.ptr<i32>
+    llvm.store %1, %3 : !llvm.ptr<i32>
+    llvm.store %0, %5 : !llvm.ptr<i32>
+    omp.target {
+      omp.parallel {
+        %8 = llvm.load %3 : !llvm.ptr<i32>
+        %9 = llvm.load %5 : !llvm.ptr<i32>
+        %10 = llvm.add %8, %9  : i32
+        llvm.store %10, %7 : !llvm.ptr<i32>
+        omp.terminator
+        }
+      omp.terminator
+    }
+    llvm.return
+  }
+}
+
+// CHECK: call void @__omp_offloading_[[DEV:.*]]_[[FIL:.*]]_omp_target_region__l[[LINE:.*]](ptr %{{.*}}, ptr %{{.*}}, ptr %{{.*}})
+// CHECK: define internal void @__omp_offloading_[[DEV]]_[[FIL]]_omp_target_region__l[[LINE]](ptr %[[ADDR_A:.*]], ptr %[[ADDR_B:.*]], ptr %[[ADDR_C:.*]])
+// CHECK:  %[[STRUCTARG:.*]] = alloca { ptr, ptr, ptr }, align 8
+// CHECK: %[[GEP1:.*]] = getelementptr { ptr, ptr, ptr }, ptr %[[STRUCTARG]], i32 0, i32 0
+// CHECK: store ptr %[[ADDR_A]], ptr %[[GEP1]], align 8
+// CHECK: %[[GEP2:.*]] = getelementptr { ptr, ptr, ptr }, ptr %[[STRUCTARG]], i32 0, i32 1
+// CHECK: store ptr %[[ADDR_B]], ptr %[[GEP2]], align 8
+// CHECK: %[[GEP3:.*]] = getelementptr { ptr, ptr, ptr }, ptr %[[STRUCTARG]], i32 0, i32 2
+// CHECK: store ptr %[[ADDR_C]], ptr %[[GEP3]], align 8
+// CHECK: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @1, i32 1, ptr @__omp_offloading_[[DEV]]_[[FIL]]_omp_target_region__l[[LINE]]..omp_par, ptr %[[STRUCTARG]])
+
+
+// CHECK: define internal void @__omp_offloading_[[DEV]]_[[FIL]]_omp_target_region__l[[LINE]]..omp_par(ptr noalias %tid.addr, ptr noalias %zero.addr, ptr %[[STRUCTARG2:.*]]) #0 {
+// CHECK: %[[GEP4:.*]] = getelementptr { ptr, ptr, ptr }, ptr %[[STRUCTARG2]], i32 0, i32 0
+// CHECK: %[[LOADGEP1:.*]] = load ptr, ptr %[[GEP4]], align 8
+// CHECK: %[[GEP5:.*]] = getelementptr { ptr, ptr, ptr }, ptr %0, i32 0, i32 1
+// CHECK: %[[LOADGEP2:.*]] = load ptr, ptr %[[GEP5]], align 8
+// CHECK: %[[GEP6:.*]] = getelementptr { ptr, ptr, ptr }, ptr %0, i32 0, i32 2
+// CHECK: %[[LOADGEP3:.*]] = load ptr, ptr %[[GEP6]], align 8
+
+// CHECK: %[[VAL_A:.*]] = load i32, ptr %[[LOADGEP1]], align 4
+// CHECK: %[[VAL_B:.*]] = load i32, ptr %[[LOADGEP2]], align 4
+// CHECK: %[[SUM:.*]] = add i32 %[[VAL_A]], %[[VAL_B]]
+// CHECK: store i32 %[[SUM]], ptr %[[LOADGEP3]], align 4
+
+// -----
+