<table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Issue</th>
<td>
<a href=https://github.com/llvm/llvm-project/issues/62389>62389</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>
OpenMP 5.1 variant function are not called within Dispatch is_device_ptr directive
</td>
</tr>
<tr>
<th>Labels</th>
<td>
</td>
</tr>
<tr>
<th>Assignees</th>
<td>
</td>
</tr>
<tr>
<th>Reporter</th>
<td>
nolanbaker31
</td>
</tr>
</table>
<pre>
When using OpenMP 5.1 directive **dispatch** with the **is_device_ptr** clause, seems to simply ignore the function call. Function add is called, and neither the device variant or the regular function are called.
```
//===---- test_dispatch_is_device_ptr.c -----------------------------------------===//
//
// OpenMP API Version 5.1
//
// Uses dispatch construct as context for variant directive. Uses is_device_ptr
// clause to ensure that the list item is in the device region.
//
//
//===-------------------------------------------------------------------------------===//
#include <omp.h>
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include "ompvv.h"
#include <stdbool.h>
#define N 4
int errors = 0;
int i = 0;
int *arr;
int host_arr[N] = {0};
void add_v(int *arr);
#pragma omp declare variant(add_v) match(construct={dispatch})
void add(int *arr){
for (int i = 0; i < N; i++){
host_arr[i] = host_arr[i] + 4; // Base function adds 4 to array values
}
}
void add_v(int *arr){
#pragma omp target is_device_ptr(arr) map(tofrom: host_arr)
for (int i = 0; i < N; i++){
host_arr[i] = arr[i] + 7; // Variant function adds 7 to array values
}
}
int test_wrapper() {
int t;
errors = 0;
t = omp_get_default_device();
arr = (int *)omp_target_alloc( sizeof(int)*N, t);
if(NULL == arr){
OMPVV_WARNING("Allocation using omp_target_alloc failed");
exit(1);
}
#pragma omp target is_device_ptr(arr)
{
#pragma omp parallel for
for(int i = 0; i < N; i++){
arr[i] = i;
}
}
#pragma omp dispatch is_device_ptr(arr)
add(arr);
for(int i = 0; i < N; i++){
printf("host_arr[%d] = %d \n",i,host_arr[i]);
if(host_arr[i] != i + 7 || host_arr[i] != 4)
errors++;
}
if(errors > 0){
printf("Dispatch w/ novariants true is not working properly");
}
OMPVV_ERROR_IF(errors > 0, "Dispatch w/ novariants true is not working properly");
return errors;
}
int main () {
test_wrapper();
}
```
Test compiles, runtime output is:
```
test_dispatch_is_device_ptr.c.o:
[OMPVV_ERROR: test_dispatch_is_device_ptr.c:58] Dispatch w/ novariants true is not working properly
[OMPVV_ERROR: test_dispatch_is_device_ptr.c:64] Condition test_wrapper() failed
[OMPVV_INFO: test_dispatch_is_device_ptr.c:63] Test is running on device.
host_arr[0] = 0
host_arr[1] = 0
host_arr[2] = 0
host_arr[3] = 0
[OMPVV_INFO: test_dispatch_is_device_ptr.c:65] The value of errors is 8.
[OMPVV_RESULT: test_dispatch_is_device_ptr.c] Test failed on the device.
```
Compilation command is:
`
clang -I./ompvv -lm -O3 -fopenmp --offload-arch=native -fopenmp-version=51 -DVERBOSE_MODE=1 /autofs/nccs-svm1_home1/nolanbaker/sollve_branches/sollve_vv/tests/5.1/dispatch/test_dispatch_is_device_ptr.c -o bin/test_dispatch_is_device_ptr.c.o 2>&1 | tee -a logs/test_dispatch_is_device_ptr.c.log && echo "PASS" > _ompvv_temp_result_.exitstatus.24964 || echo "FAIL" > _ompvv_temp_result_.exitstatus.24964`
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJysWFuPm7oW_jXOyxIoGAjMQx5ymRyN1Gaq6e7sR-SAE3y2sZFt0t3z649sLoFkMm1HRSgB2-v2rZsN0ZqdBKVLFK9RvJ2RxpRSLYXkRBzIP1SFwewgix_Lv0sqoNFMnOC5puLzF4j9AAqmaG7YmQLCK4RXBdM1MXnZvsF3ZkowZT_LdFbQM8tpVhvVLck5aTRFeAOa0kqDkaBZVfMfwE5CKurIj43IDZMCcsK5D7v-lRQFMO1GaWF5EFGAoMyUVDnCVhyciWJEGJDtqKKnhhN1YUsU7Zj4aL5F81X3u5h3d_uKd_YOt-3teZ4HhmqT9VZnEwP9HLxfvQamnYyxvNFzD_3qyxO8UqWt7rEf3F3-TVMNvXaQS6GNanIDRNsXQ_81cJRqgGfwpt9STv01Ztx6zTqLCt04LxHjsOVMG2CGVtYxTIy9oOiJSeHf0_YuyH_wuoNzyETOm4ICCjeyqv0ShY9vzWlTMPneLGeHu9MVMeWdSYxlVZ_PfokwvsP5ICWfUhf0yASFPUTtEBMGqFJSaUDhFuYoXF8m2BtjCK-IUpOhUmqT2cF4vUfx1hGhZD1HyXZYd5assJmXnRFOR3zww7AE4bBW5FQRkFUNBc25zbAuzBBOO-oHqNpikQ6haV2TrIcykmwt24nYG6FJJxRcLHfTF3vd4wb27hHhtbvHVDC2mvVW34zhNUSWR5cBa6JHdYkUhYbI5gNRivyAM-EN1b0Ea0aLS__wDoYXxaYoGqJO1FzlJE5bKqhIjXBq5FHJCoWri_o9fH8WnStgkhEwr10xmWKT_Do2VkFXVr8rUtfUGmktHOnkVgzBBm9HPRg3IKs6O1GTFfRIGm468FqeIxZEqTbYB3cg_GBpW9gzwrnMEU5Bs_9ReWyXWRZ4tbeNx0zZMbti_-3TJ2jLDdw4117Pn7-8vmZ_r172T_v_OJXwysohDre22V7rAEfCXLPDU4n2ov8ym13BdGYA9rci6kI9EjGlr4myPZPbyOrW2KcPhBjAVXyxiWVjC-4ZM3S5982BroRcFSz4uOa1YsIcW--NUgXhuBgqKI4LQPFGOK9tGMKbq6Sa6OKC57b-BA6YNt8AJRuUbN6oUm5VNLG4Cw6XJJ0Vg6wRnk7skEqPML-2dWzptof7u815IbvarsGohtrWL6SB71L9Y2O4VrKmiv-4CdqR9DYXHl9enl-yp92NJhv4Y0IVNY0SPR5Dyxrq0FCEKsIE3FSf29L0Bo-rjaP7_YtqA7msasapthapRhhWUZCNqRubhyhcvbP9fHer6csLcbwegWmbwbuUKFzFqY2dj6H7QZGLyIqEjRQFc9XujYrfVbqphKf97vmXBIRWgMOcaQu1cNVUdJvRbhc6yp95n63zm6ng_hS-PxVeTf2-CbEzoaRtywR57Dsd05D6V1xfHr9--_TXz_n2qLToWkQuW3T_nfDduMhtW1Muq8oetcYh26_OOREn8J58hHduVwser8B7DsE7ypqKqgbPk8cjl6TwiMpLFG4FcWfIfoF3bs82KNzGAXjb18eX9fPXx-zz8_YRhdvAbjVIY-RRI7wTea49fa6CrJQVDezIcHhFeKcl52eaHRQReWnzrh85nxHeWajsmD1E4d3l9Lr7yblOwoGJny3zJWC7YceLwNZrMJSCR4DLk_4pKZcnQHiB8AJoXkpb_b6svn5FGLuKmDlkM0OrOlNU242Nb7u_NsQ02sfRwyLqe0RPv1s9ffoderSYz4plWDyED2RGl8EixXGYpPNoVi7ToIjSiAZhSnBIj0kexQkOyTGIiiSc5_GMLfEch_MIL4KHcBFF_iKlJDngJDjQhORBiqI5rQjjPufnypfqNGNaN3S5wGH6MOPkQLnuP0qopV3kHZqTRtHcnjD1hcwww-ly9E3ifLMBVdTVrvaE7z5KMHEpdxPoL2fgWaP4sjSmdjHu9rYnZsrm4OeyQnhnFej-vFrJ_9LcILxzRlj3Ojv-HwAA___VjiHb">