<table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Issue</th>
<td>
<a href=https://github.com/llvm/llvm-project/issues/63555>63555</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>
[MLIR][OpenMP] Optimizations/analyses cross omp.target op region boundaries.
</td>
</tr>
<tr>
<th>Labels</th>
<td>
openmp,
mlir
</td>
</tr>
<tr>
<th>Assignees</th>
<td>
jsjodin
</td>
</tr>
<tr>
<th>Reporter</th>
<td>
jsjodin
</td>
</tr>
</table>
<pre>
Compiling OpenMP with target offloading is done in two passes. The code is compiled for the host, then for the target device. When compiling for the device the code outside a target region (inside a host function) cannot be used to optimize the code inside the target region. The reason for this is that the environment is different compared to the host e.g. macro values may be different.
The inputs and outputs to the target region cannot change because moving code across a region boundary may cause incorrect code.
There is a discussion regarding this on discourse:
https://discourse.llvm.org/t/rfc-prevent-optimization-analysis-across-omp-target-region-boundaries/71402
For example if the host and device had different sizes for "int"
int max0 = INT_MAX; // 32-bits
#pragma omp target
{
int max1 = INT_MAX; // 16-bits
if (x == max1) {
.. do something..
}
}
In this case if max1 was replaced with max0 and passed in to the target region it would get the incorrect value.
There is currently no way to have a barrier that prevents optimizations and analyses to propagate information across target op region boundaries. For now the proposed solution when doing device compilation is to outline the target ops into separate functions that only contain the target op (and a return). This should prevent any illegal transforms from happening, but if some transform still moves code outside the target region it is easily detected since only 2 ops should be present in the outlined function.
Small test, run with "mlir-opt -cse" and c1 will be replaced with a constant outside the omp.target op.
```
func.func @constant_hoisting_target(%x : !llvm.ptr<i32>, %y : i32) {
%c0 = arith.constant 10 : i32
%z = arith.addi %y, %c0 : i32
omp.target {
%c1 = arith.constant 10 : i32
%w = arith.addi %z, %c1 : i32
llvm.store %w, %x : i32, !llvm.ptr<i32>
omp.terminator
}
return
}
```
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJyMVt1u4zYTfRr6ZmDBpi3_XPhik3wGFvi2W7QLtHeLMTWWuKBIgaTsOE9fDEUp9m4KNAgSSZy_M3PmSBiCri3RQZRPQsof4YertBVSivJlhn1snD_kh7OTq26HZ9d22mhbw9eO7Jff4apjAxF9TRHc-WwcVnyqA1TOEmgL8eqgwxAoFPCtIVCuIj5XKRRVcHYeYkPQuBCFfOZrOz3MoSu6aEUF_MVnaipitBqO02UK7_oYdEWAo7-nWjsLQu60zSecDs69VVE7K-QeFFrrIpwI-kAVRAeui7rVb3eBs_ddZUPkAZonDG6sXQdGGRuMyZzsRXtnW7IxdUefz-T5htGgHxKObQAq6gJaVN7BBU1PAVq8cWmTXwFi8SIWnzivtl0fA6CtGHm6zsEe4WeEqkFbE5xIYR8IWnfhXiZ8nDEEwNHj5Hpbob-l9IO5tsp5Tyomj2KoYqrFp9kiQKWD6kPgIJ5q9IkWqSvOpkPX-0BilV2bGLvAd_Io5HE6L4y5tIXztZDHKOTRn9W883QhG-d5Osjzm6NFcws6zAcEc9d28wH8fIAyz1A0BSGP2-V6Ie9LPzoP9IptZwj0-X0Q3NPMrgaru7EF_UYhjVpIqW3kpUmRtI3Q4usCxOoFPv_27fuXT3-L1RMMyGAl5ycdQ84tV53HukVwbTcOKw9WbJ-GC4AccvkvIZebu5AAwACE3L2yOXuwKxP8LiL_FAVUDoJrKTba1kUeJYjty1hBvvhsh8kpDKk7qZYrBvDUGVRUDTKQUHPD0rpXafk_oqGOcHW9qYCfxeaeUonsH3NK9Z4bb25gHVzxxrEbvPAqn9B7TX5YtkyPAPf8GJZjIAml7ei867DGyNnPzrfJbOT_KGjdT3ugWcSAuWLdNZXOYRyDDc70KcSVJapyTPdMnEGvhgQ6JXd9NNo-6IjrAs_ZQaAOPdc1alMWEWfNDZSzEbmv94487AQPPMXes5ixHukAoUl9zi0BtDfQxlCNBqJHGxh4gLN3LTTYdWS1rVmCT33kOTM33g0hRG0MywWFR5n9cMQ6AGHQ5gYVRVKRe6StogGITIBzfSfuI4UkjQO23KBqasIDJ_5s0RiINLwvfG8H_gkpW6M9CwPMVSAhZRq7WsKVKz_RT4RF7meIaOMDFNd2xdTch8Ris8i_6ZaLK_gPiPViDPW9cTpEbevvQwwhd0KWvIyfQMhlErQuerF61ispVv9jCEKWt2TAjx43VchSDVqCXsemmCpeLiaPd9O3O0usKp0i5wzqF4c7oA_awMbL_5YzGV8_yPo2Zl3-6pJ6EKLzlLyz5et7B54_7tQUIBVOvtUWo_MT_lGvYFyERxmbpjerDqtqv9rjjA7LzW672-5LuZ01h9Vis92td4v1tlS4IbWX6_VGlee13G-XuNnO9EEu5GqxkdtlKUu5Lfbb5Xa3wpMslSxPy71YL6hFbaYX10yH0NNhsyrLcmbwRCbkry3XkW07fm_I50zd_OXlD-w9P_V1EOuF0SGG93hRR5M-2L78__MfonwR5dPwLSbKF_h6r3lCHifFG3TtgdkfaNus9-bw-CqudWz6U6FcK-SRi8j_5p13P0jxSzkh5GwJ5D8BAAD__99jY9g">