| Issue |
146280
|
| Summary |
clang does not compile is_device_ptr when the pointer is a member of a struct...
|
| Labels |
clang
|
| Assignees |
|
| Reporter |
bschulz81
|
If an algorithm becomes more complex than a very simple case, it often needs to store temporary data.
Say you want to create a tensor struct, which has data, extents and strides for multi dimensional arrays.
Then, the natural course would be to use omp_target_alloc, to allocate some space, let the pointers, which are members of the struct point to that, and then do a loop. As the members are device pointers, you would want to put in something like
is_device_ptr(arraystruct.devptr)
Unfortunately, is_device_ptr does not follow the dots and stops compilation with
/openmptestnew/openmpoffloatest/main.cpp
/home/benni/projects/openmptestnew/openmpoffloatest/main.cpp:16:36: error: expected variable name
16 | #pragma omp target is_device_ptr(u.t)
. Note that this is unlike the map command.
Here is a non-working example where clang stops compilation.
```
#include <omp.h>
struct arraystruct
{
int A;
double* t;
};
int main()
{
arraystruct u;
u.t=(double*)omp_target_alloc(sizeof(double)*20,omp_get_default_device());
#pragma omp target is_device_ptr(u.t)
#pragma omp teams distribute
for(size_t i=1;i<20;i++)
{
u.t[i]=20;
}
double *g=u.t;
}
```
note that we could get u.t still onto the device by using map as follows
```
#include <omp.h>
struct arraystruct
{
int A;
double* t;
};
int main()
{
arraystruct u;
u.t =new double[20];
#pragma omp target enter data map(alloc:u.t[0:20])
#pragma omp target teams distribute
for(int i=1;i<20;i++)
{
u.t[i]=20;
}
```
The problem with that version is just that for map, you could always map down, and you have to associate the pointer. When the association is made, it is checked whether memory of the same size was allocated by the host.
Say your large language model needs a temporary data storage of 20 gigabyte on vram...
It is just absolutely annoying, when you then have to allocate this on disk too, even if you absolutely do not need it because you just want a temporary data space on device, and that is_device_ptr does not compile if your pointer is a struct member..
This behavior is wrong, according to the standard. Is device_ptr should accept any C-type pointer.
In the example above, double *g=u.t; does not need any conversion, so u.t is definitely of type C-type pointer. It is just that is_device_ptr does not accept it as the compiler refuses to follow the dots....
![]()
_______________________________________________
llvm-bugs mailing list
llvm-bugs@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-bugs
