[llvm-bugs] [Bug 50968] New: target amdgcn-amd-amdhsa: accessing struct member array >64bytes yields [GPU Memory Error]

[via llvm-bugs]

https://bugs.llvm.org/show_bug.cgi?id=50968

            Bug ID: 50968
           Summary: target amdgcn-amd-amdhsa: accessing struct member
                    array >64bytes yields [GPU Memory Error]
           Product: clang
           Version: trunk
          Hardware: PC
                OS: Linux
            Status: NEW
          Severity: enhancement
          Priority: P
         Component: OpenMP
          Assignee: unassignedclangbugs at nondot.org
          Reporter: gcc.j.kelling at hzdr.de
                CC: llvm-bugs at lists.llvm.org

Compiling the following code:
```c++
#include <iostream>

template<int N>
struct A
{
        char arr[N];
        int b = 0;

        A()
        {
                arr[0] = 0;
        }
};

template<int N>
void f()
{
        int b = 0;

#pragma omp target map(tofrom: b)
        {
#pragma omp teams distribute
                for(int i = 0; i < 1; ++i)
                {
                        A<N> a;
                        b = a.arr[0]; // (1) only this fails
                        b = a.b += sizeof(a); // (2)
                }
        }

        std::cout << b << std::endl;
}

int main()
{
        //OK:
        f<64>();

        //[GPU Memory Error] Addr: 0x0 Reason: Page not present or supervisor
privilege.:
        f<65>();
}
```

with
clang++ -fopenmp -fopenmp=libomp -fopenmp-targets=amdgcn-amd-amdhsa
-Xopenmp-target=amdgcn-amd-amdhsa -march=gfx900 memberArray.cpp

and running on Vega 10, yields:
```
[GPU Memory Error] Addr: 0x0 Reason: Page not present or supervisor privilege. 
Memory access fault by GPU node-6 (Agent handle: 0xbd2010) on address (nil).
Reason: Page not present or supervisor privilege.
Aborted (core dumped)
```

Note, that this happens only when the array access in line (1) is present, the
access to the other member in line (2) does not cause a memory error and the
value returned by `sizeof(a)` is as expected. Also, this happens if the member
array is larger than 64 bytes (any type, e.g. int[16] works, int[17] does not).

Expected behavior: No size restriction on (fixed-size) member arrays.

Rationale: We require a small-object allocator in block
shared(NVIDIA)/local(AMD) memory on GPU, for which we require a type with a
member array filling up most of the available shared memory  (>~30kB).

-- 
You are receiving this mail because:
You are on the CC list for the bug.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-bugs/attachments/20210702/df347951/attachment.html>