<table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Issue</th>
<td>
<a href=https://github.com/llvm/llvm-project/issues/76577>76577</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>
[MLIR][OpenMP][Libomptarget] `omp.target` access memory raise errors
</td>
</tr>
<tr>
<th>Labels</th>
<td>
mlir
</td>
</tr>
<tr>
<th>Assignees</th>
<td>
</td>
</tr>
<tr>
<th>Reporter</th>
<td>
EllisLambda
</td>
</tr>
</table>
<pre>
The MLIR and LLVM toolchain was built with https://github.com/llvm/llvm-project/commit/8c6172b0ac2b254dec7d57326abfd666a7954a03. The MLIR code use `mlir-translate --mlir-to-llvmir| clang++ -c -x ir -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx1103` to generate static lib
```
module attributes {llvm.target_triple = "amdgcn-amd-amdhsa", omp.is_gpu = true, omp.target = #omp.target<target_cpu = "gfx1103", target_features = "">} {
llvm.func @llvm_omp_target_alloc_device(i64, i32) -> !llvm.ptr
llvm.func @omp_get_default_device() -> i32
llvm.func @_QQmain_omp_outline_1() attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost)>} {
%0 = llvm.mlir.zero : !llvm.ptr
%1 = llvm.call @omp_get_default_device() : () -> i32
%2 = llvm.getelementptr %0[67108864] : (!llvm.ptr) -> !llvm.ptr, f64
%3 = llvm.ptrtoint %2 : !llvm.ptr to i64
%4 = llvm.call @llvm_omp_target_alloc_device(%3, %1) : (i64, i32) -> !llvm.ptr
%5 = llvm.call @llvm_omp_target_alloc_device(%3, %1) : (i64, i32) -> !llvm.ptr
%6 = omp.map_info var_ptr(%4 : !llvm.ptr, f64) map_clauses(tofrom) capture(ByCopy) -> !llvm.ptr
%7 = omp.map_info var_ptr(%5 : !llvm.ptr, f64) map_clauses(tofrom) capture(ByCopy) -> !llvm.ptr
omp.target map_entries(%6 -> %arg0, %7 -> %arg1 : !llvm.ptr, !llvm.ptr) {
^bb0(%arg0: !llvm.ptr, %arg1: !llvm.ptr):
%8 = llvm.mlir.constant(0 : index) : i64
%9 = llvm.mlir.constant(1 : index) : i64
%10 = llvm.mlir.constant(8192 : index) : i64
omp.teams {
omp.parallel {
omp.wsloop for (%arg2, %arg3, %arg4, %arg5) : i64 = (%8, %8, %8, %8) to (%10, %10, %10, %10) step (%9, %9, %9, %9) {
%11 = llvm.mul %arg2, %10 : i64
%12 = llvm.add %11, %arg3 : i64
%13 = llvm.load %arg0 {alignment = 8 : i64} : !llvm.ptr -> vector<16xf64>
%14 = llvm.load %arg1 {alignment = 8 : i64} : !llvm.ptr -> vector<16xf64>
%15 = llvm.mul %arg3, %10 : i64
%16 = llvm.add %15, %arg2 : i64
%17 = llvm.getelementptr %arg1[%16] : (!llvm.ptr, i64) -> !llvm.ptr, f64
%18 = llvm.load %17 {alignment = 8 : i64} : !llvm.ptr -> vector<16xf64>
%19 = llvm.fmul %13, %14 : vector<16xf64>
%20 = llvm.fdiv %14, %18 : vector<16xf64>
%21 = llvm.fadd %19, %20 : vector<16xf64>
%22 = llvm.getelementptr %arg1[%12] : (!llvm.ptr, i64) -> !llvm.ptr, f64
llvm.store %21, %22 {alignment = 8 : i64} : vector<16xf64>, !llvm.ptr
omp.terminator
}
omp.terminator
}
omp.terminator
}
omp.terminator
}
omp.barrier
llvm.return
}
llvm.func @_mlir_ciface__QQmain_omp_outline_1() attributes {llvm.emit_c_interface} {
llvm.call @_QQmain_omp_outline_1() : () -> ()
llvm.return
}
}
```
Using C program to call the function and build with clang args `-fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx1103 ` it's normal when running on the CPU host, but on AMDGPU it raise errors even if replace `llvm_omp_target_alloc_device` to `llvm.alloc` and with a smaller size:
```
Libomptarget error: Host ptr 0x0000560ed81500a1 does not have a matching target pointer.
Libomptarget error: Consult https://openmp.llvm.org/design/Runtimes.html for debugging options.
Libomptarget error: Source location information not present. Compile with -g or -gline-tables-only.
Libomptarget fatal error 1: failure of target construct while offloading is mandatory
```
The C code has similar function works normally on the AMDGPU:
```
int main() {
double* a = (double*)llvm_omp_target_alloc_device(4096*4096*sizeof(double) , omp_get_default_device());
double* b = (double*)llvm_omp_target_alloc_device(4096*4096*sizeof(double) , omp_get_default_device());
#pragma omp target teams map(tofrom: a, b)
#pragma omp parallel for
for(int i=0; i<4096; i++){
for(int j=0; j<4096; j++){
for (int k=0; k<4096; k++){
a[i*4096+j] = i * j;
b[j*4096+k] = j / i;
}
}
}
}
```
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzMWN1u67gRfhr6ZmBDoqy_C1_EdtIukKBnT_cUvTMoiZKZUKRAUvnZpy9ISZbsKHFOi4OukcA0xflm5uNwOCOiNasEpRsUblG4X5DWHKXa3HLO9D2ps4IsMlm8bf44Uni4_-07EFHA_f2_HsBIyfMjYQJeiIasZdzACzNHOBrTaBTcIHyH8F3FzLHNVrmsEb7j_Hn4WjZKPtLcIHyXy7pmdpDkkR_jzCM5znC4LmgeF2Ec4IhkZRFFEYnTcE28YAUnc3JZUGg1BRR5NWdqaRQRmhNDYbnsJuTS6mMKxTvIOREVwluEt7DMYfkKTMGylA0VdXMaLA1RFTUaBXtSF1UulqQu7P9RE1j--2zR_JqaqPyIgn1Vvvq-F6DIAyOhooIqa5o2xLAcOMuQt0feDYq8_s_9rGXRcgrEGMWy1lANKN5aJ1adzoNRrOEUULAHhPE7_QhjhHcg62bF9KFqWrfSqJYO0x1ODxCMMyjY9SryXgphPDjRofbPS0pMq6xp3SL7F9yieG9t7dwAZ3LZihzQ2rM_DrJuDj0A4Vzmh4I-s5winLBobdFZgBFOYYmCW0DYdwiNUT3gOaIFs0gFLUnLzYh1QrBoc5KH33-vCRPOHNkazgQ9-L3gOeuWmoLmnCh6eE9a_-TEXWfBwbw1w-YkQh6lNgin79gBAIRDzy101tlwXf1JlQQU3Mx57wT8USAnnF8losOa5cTh4RGvooZyWlNhGqOccSjcRrHvJUm0RuF-xBpNm9ssvIMyWp9pCUYtjVFGMmEG5Weu2mPCLmTX7z2-EkxWobXC0jXh4Gsx5nSG_wedkdNp46omzYGJUsIzUQfHaNLzcDNHNE7BSuSctJpqhBMjS2UTbgo5aewxRTjZvu1k83bViPiKEeEvN2KSnywgFUYxB-go6uVCoiqvpzueTvpz9l0E7HgGUXibZV6H7RDnZB3suyepveJORjvukouznEuhDREG4cRzZjFR0NchOM7C3MmnH8v7X5H3L5PJBCDxU3wVw1FPSa0nFNm5hijCOeXn2WuUedFcygZKqQBOZOKRvWAcrsdhOLFiSJeWxX7FzHdq00O3yh92f36Qgja06dem_YOZ77NgwKE_ya51y-HcE9-DOdZG_ie5lBRFBzhh4Yr0JEdySYohyq2FhLNK2MTsliQnnHj_LoG6s_BMcyMVCnZ-9GoPZ3A7r3I9q9L_lSrDOYaDrzIcvWc4HBnGV6TjT-46d8jDrVPy0V23c9BXrzyLkbxn1qr_dbxOkkfZE-ufaF13xHyIZa_iSfIoC_bcyQ0AyRWAc2Pw5BiVwz4Npw57P4f1WYEy2TT8P27a5OMWaCMV7ZwZDMdf2L8Zry4uoDGtGqpqJohdf2EDivdzefYjgXH5x6suID9aeA6VEaUYnTx2XihqWiWGyQnweZFtr6BDzkqS08PXK26HQWtmDvmBCUOVlX9XOZ9VZp-BX5a_3fgzh07-jIPz9uyHZqKCHTRKVorU9k5ylpgjBeu8YVK4Ptk2xUXXFLvGE4iqtO1Uv9ZontrS_6LjtFrANtWxBiFVTTi8HKkA1QphrZfCmbv79gO6_mQHWWvs9M3D_m_ffgAzoAjTFKhSUmmgz1QAK0HRhpPc9dufFsRdx9uvWrlHds7S4gghoGtbUyjQ7E96qqUuqL5nmaybvhp0ltj9_LvUBmwG8F49z_PCyKNF4oeeR3woJLUeGziSZwoEamLyo_W4B2mkC6rVZ_g7KXTLzcWbjG4TVs4fqSqE7wqqWSUQvvveCsNqqldHU3NXBhU0a6vKMd3YeNCfKvynbFVOgcucuOCxZbequ7H1pVFUU2FWsJN1wzjtKFxWIBUsKxvvS0MyTvVSCv42p6okhvBOIbhatiSMt4qCLAdmXLGo2tzAy9HqkGVpby7rA9NQE1HYTPE2u09_2FjqXsYciQbNasaJGk_Di1RPQyDytyH6ulD7aO9ti2iP9XCQT8e_kG3GKcI3QIay8TSFcHqlT1t7aYTwTf9lg0-WE4QU-nckH3bUrvSfpKLRnOyvYA7CQaNIVRMrNextV9TXpBn7suAGiDv2Z-nwXPpU95eXN46dwIndIoaCvYeCrR3snBtu7N6xWejTto0ij4PI40TkcUbk4mMPVg_xNEA8TSCerkMQFG7Zie_tY1c07IGB3cDHMyannwyF28dR7mmQewSE76zrH8jNXOQzk59cOYtiExRpkJIF3fixFyRxmOBocdx4Se4VuAwDkqW-F_hR6ZM8j-i6WPsF9ZMF22APBz7GqZ-EHo5WcVp46yxN1sk6pUXuobVHa8L4KaUtmNYt3cRRGMcLTjLKtXsvjLG9yhG2JdZCbdzr26ytNFp7nGmjRwDDDHfvkh_uf_uOwj0Kt_9oqHj41o2nKckxGHmTd4_2eshzqjXUtJbq7ewCWrSKb3761bJzSCN853z6TwAAAP__gw-RKQ">