when offloading nested structs, sometimes "PluginInterface" error: Failure to synchronize stream (nil): Error in cuStreamSynchronize: an illegal memory access was encountered happens

[bschulz81]

The following code attached to this bugreport has no memory problems. but it fails with

Entirely on gpu
"PluginInterface" error: Failure to synchronize stream (nil): Error in cuStreamSynchronize: an illegal memory access was encountered
omptarget error: Consult https://openmp.llvm.org/design/Runtimes.html for debugging options.
mdspan_acc.h:3807:9: omptarget fatal error 1: failure of target construct while offloading is mandatory

Process returned -1 (0xFFFFFFFF) execution time : 0.394 s

Upon investigating, it seems not to offload dQ and dR correctly in the function qr_decomposition in line 3792 which calls
create_in_struct(dA);
create_out_struct(dQ);
create_out_struct(dR);

that do the mappings with

template
void inline create_in_struct(const datastruct& dA)
{
#pragma omp target enter data map(to: dA,dA.pdata[0:dA.pdatalength],dA.pextents[0:dA.prank],dA.pstrides[0:dA.prank])

}

template
void inline create_out_struct(datastruct& dA)
{
#pragma omp target enter data map(to: dA) map(alloc: dA.pdata[0:dA.pdatalength]) map(to:dA.pextents[0:dA.prank],dA.pstrides[0:dA.prank])
}
template

and then calls

gpu_qr_decomposition.

If one removes the lines 1863 -1997 in gpu_qr_decomposition, especially these lines:

#pragma omp parallel for
for (size_t i=0; i<Q.pdatalength; i++)
{
Q.pdata[i]=0;
}
//
//
#pragma omp parallel for
for (size_t i=0; i<R.pdatalength; i++)
{
R.pdata[i]=0;
}

then the code suddenly compiles...

The strange problem is that before, similar code is called for a cholesky and an lu decomposition.

here, the code works provided i compile it without optimization.

if i compile the code with -O2, then clang takes veeeeeeerry long to finish, and for the result, the lu decomposition crashes too...

There is no problem with the array sizes, as I have checked these, and these are just test cases with 9 elements (3x3 matrix).

main_acc.cpp.txt

mdspan_acc.h.txt

CMakeLists.txt

[bschulz81]

clang++ --version
clang version 20.0.0git97c3a990
Target: x86_64-pc-linux-gnu
Thread model: posix
InstalledDir: /usr/lib/llvm/20/bin
Configuration file: /etc/clang/x86_64-pc-linux-gnu-clang++.cfg

[llvmbot]

@llvm/issue-subscribers-offload

Author: None (bschulz81)

The following code attached to this bugreport has no memory problems. but it fails with

Entirely on gpu
"PluginInterface" error: Failure to synchronize stream (nil): Error in cuStreamSynchronize: an illegal memory access was encountered
omptarget error: Consult https://openmp.llvm.org/design/Runtimes.html for debugging options.
mdspan_acc.h:3807:9: omptarget fatal error 1: failure of target construct while offloading is mandatory

Process returned -1 (0xFFFFFFFF) execution time : 0.394 s

Upon investigating, it seems not to offload dQ and dR correctly in the function qr_decomposition in line 3792 which calls
create_in_struct(dA);
create_out_struct(dQ);
create_out_struct(dR);

that do the mappings with

template<typename T>
void inline create_in_struct(const datastruct<T>& dA)
{
#pragma omp target enter data map(to: dA,dA.pdata[0:dA.pdatalength],dA.pextents[0:dA.prank],dA.pstrides[0:dA.prank])

}

template<typename T>
void inline create_out_struct(datastruct<T>& dA)
{
#pragma omp target enter data map(to: dA) map(alloc: dA.pdata[0:dA.pdatalength]) map(to:dA.pextents[0:dA.prank],dA.pstrides[0:dA.prank])
}
template<typename T>

and then calls

gpu_qr_decomposition.

If one removes the lines 1863 -1997 in gpu_qr_decomposition, especially these lines:

#pragma omp parallel for
for (size_t i=0; i<Q.pdatalength; i++)
{
Q.pdata[i]=0;
}
//
//
#pragma omp parallel for
for (size_t i=0; i<R.pdatalength; i++)
{
R.pdata[i]=0;
}

then the code suddenly compiles...

The strange problem is that before, similar code is called for a cholesky and an lu decomposition.

here, the code works provided i compile it without optimization.

if i compile the code with -O2, then clang takes veeeeeeerry long to finish, and for the result, the lu decomposition crashes too...

There is no problem with the array sizes, as I have checked these, and these are just test cases with 9 elements (3x3 matrix).

main_acc.cpp.txt

mdspan_acc.h.txt

CMakeLists.txt

[jhuber6]

Can you format the issue to make it easier to read? If this gives the expected results at O0 but fails with higher levels of optimization it might just be a compiler bug, but it's difficult to tell with such a large reproducer. This seems to only be a single file so I wouldn't expect LTO to make difference here.

I can't really do much with what you've given me, seems to depend on openacc.h. If you think it's a compilation bug I would recommend compiling it something like this and increasing <N> until you get the pass that introduces the failure first. After that you can try to get a reproducer by checking the output temp bitcode that we can then run the pass manually on via opt.

$ clang++ bug.cpp -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -Xarch_device -mllvm=-opt-bisect-limit=<N> -O2 -save-temps

At some point I should probably write some documentation on creating reproducers for bugs like this, the process is not intuitive. It's uncommon for bugs to get fixed when they live inside 100k+ lines of LLVM-IR.

[bschulz81]

Hi the openacc.h is a leftover which is there since i converted it from openacc to openmp. it can be removed safely.

[bschulz81]

the code makes a matrix multiplication, then a cholesky and then a lu and then a qr decomposition. without O the qr decomposition fails, with O2 the lu decomposition fails with memory errors.

[bschulz81]

here are the files without openacc.h.. i.e. with pure openmp code.

main_acc.cpp.txt

mdspan_acc.h.txt

CMakeLists.txt

the output without o is:

output-without-o.txt

Note that the qr decomposition gets a memory exception because it is not offloaded correctly. but i do not think my code is at fault, because for the lu decomposition, i have to offload matrices in exactly the same way.

With -o1, this is the output

output-with-o1.txt

as you see, now the cholesky decomposition already fails

[bschulz81]

I modified the cmakelists.txt to include your debug arguments....

clang++ ./main_acc.cpp -std=c++20 -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -Xarch_device -mllvm=-opt-bisect-limit=2 -save-temps

note that it has to be linked with these (openmpi):
target_link_libraries(arraytest PRIVATE rt m c stdc++ gomp mpi)

that is the logfile which i then got without the optimization, where only the qr decomposition fails:

log.txt

that is the logfile i get at O1, where already the lu decomposition fails:

debuglog_with-O1.txt

here is the code again and the output.

Strangely, the simple matrix multiplication works. the more complex cases dont...

CMakeLists.txt

main_acc.cpp.txt

mdspan_acc.h.txt

output-without-o.txt

output-with-o1.txt

[bschulz81]

One problem when mapping this is of course that, well what it needs to map correctly is a struct of the following definition:

#pragma omp begin declare target
template
struct datastruct
{
T* __restrict pdata = nullptr;
size_t* __restrict pextents = nullptr;
size_t* __restrict pstrides = nullptr;
size_t pdatalength = 0;
size_t prank = 0;
bool prowmajor = true;

// Constructors
datastruct(
    T* data,
    size_t pdatalength,
    bool rowm,
    size_t rank,
    size_t* __restrict extents = nullptr,
    size_t* __restrict strides = nullptr,
    bool pcompute_datalength = false,
    bool compute_strides_from_extents = false
);

datastruct(
    T* __restrict data,
    size_t datalength,
    bool rowm,
    size_t rows,
    size_t cols,
    size_t* __restrict extents = nullptr,
    size_t* __restrict strides = nullptr,
    bool compute_datalength = true,
    bool compute_strides_from_extents = false
);

datastruct(
    T* __restrict data,
    size_t datalength,
    bool rowm,
    bool rowvector,
    size_t length,
    size_t* __restrict extents = nullptr,
    size_t* __restrict strides = nullptr,
    bool compute_datalength = true,
    bool compute_strides_from_extents = true
);

~datastruct();

// Operator overloads
T& operator()(const size_t* __restrict indices) __attribute__((always_inline));
T operator()(const size_t* __restrict indices) const __attribute__((always_inline));

T& operator()(const size_t row, const size_t col, const size_t strides0, const size_t strides1) __attribute__((always_inline));
T operator()(const size_t row, const size_t col, const size_t strides0, const size_t strides1) const __attribute__((always_inline));

T& operator()(const size_t row, const size_t stride) __attribute__((always_inline));
T operator()(const size_t row, const size_t stride) const __attribute__((always_inline));

// Member functions
datastruct<T> substruct(const size_t* __restrict poffsets, const size_t* __restrict psub_extents, size_t* __restrict psub_strides = nullptr, T* __restrict sub_data = nullptr);
datastruct<T>   subspanmatrix( const size_t row, const size_t col,const  size_t tile_rows,const  size_t tile_cols,  size_t * __restrict sub_extents,  size_t *__restrict sub_strides=nullptr, T* __restrict sub_data=nullptr);

datastruct<T> transpose(size_t* __restrict newextents, size_t* __restrict newstrides);
datastruct<T> row(const size_t row_index, size_t* __restrict newextents, size_t* __restrict newstrides);
datastruct<T> column(const size_t col_index, size_t* __restrict newextents, size_t* __restrict newstrides);

};
#pragma omp end declare target

this has methods, and, some variables indicating the lengths, and well 3 pointers to arrays with the lengths indicated by the member variables pdatalength for pdata and prank for the arrays pstrides and pextents. the challenge is then to map the arrays not only to the device, but also, well to map the pointers in the structsuch that they point to the arrays on the device....

The strange thing is that these functions:

template
inline void update_device(datastruct& dL)
{
#pragma omp target update to(dL.pdata[0:dL.pdatalength])
}

template
inline void update_host(datastruct& dL)
{
#pragma omp target update from(dL.pdata[0:dL.pdatalength])
}

template
void inline create_in_struct(const datastruct& dA)
{
#pragma omp target enter data map(to: dA,dA.pdata[0:dA.pdatalength],dA.pextents[0:dA.prank],dA.pstrides[0:dA.prank])

}

template
void inline create_out_struct(datastruct& dA)
{
#pragma omp target enter data map(to: dA) map(alloc: dA.pdata[0:dA.pdatalength]) map(to:dA.pextents[0:dA.prank],dA.pstrides[0:dA.prank])
}
template
inline void exit_struct(datastruct &dA)
{
#pragma omp target exit data map(delete: dA.pstrides[0:dA.prank],dA.pextents[0:dA.prank],dA.pdata[0:dA.pdatalength],dA)
}

template
inline void exit_struct(const datastruct &dA)
{
#pragma omp target exit data map(delete: dA.pstrides[0:dA.prank],dA.pextents[0:dA.prank],dA.pdata[0:dA.pdatalength],dA)
}

do that job for the matrix multiplication. yet for the qr decomposition, after offload, a simple command like

datastruct dA=A.pdatastruct;
datastruct dQ=Q.pdatastruct;
datastruct dR=R.pdatastruct;
T* __restrict buffer=(T*) omp_target_alloc(sizeof(T)(dA.pextents[1]dA.pextents[1]+2dA.pextents[0] dA.pextents[1]),omp_get_default_device());

    create_in_struct(dA);
    create_out_struct(dQ);
    create_out_struct(dR);

    #pragma omp target map(to:step_size) is_device_ptr(buffer)
    {
        gpu_qr_decomposition(dA,dQ,dR, (T*)buffer,step_size);
    }

and that here
const size_t n = A.pextents[0];
const size_t m = A.pextents[1];

if(step_size==0)
    step_size=(size_t)pow(A.pextents[0],0.8385);

const size_t nm=n*m;
T* __restrict tempC;
T* __restrict tempS;
T* __restrict tempM;

if(buffer==nullptr)
{
    tempC=(T*)omp_alloc(sizeof(T)*m*m,omp_large_cap_mem_alloc);
    tempS=(T*)omp_alloc(sizeof(T)*nm,omp_large_cap_mem_alloc);
    tempM=(T*)omp_alloc(sizeof(T)*nm,omp_large_cap_mem_alloc);
}
else
{
    tempC=buffer;
    tempS=buffer+m*m;
    tempM=tempS+nm;
}

size_t Mext[2];
Mext[0]= A.pextents[1];
Mext[1]=A.pextents[0];
size_t Mstrides[2];
Mstrides[0]= A.pstrides[0];
Mstrides[1]=A.pstrides[1];



#pragma omp parallel for
for (size_t i=0; i<nm; i++)
{
    tempM[i]=A.pdata[i];
}


#pragma omp parallel for
for (size_t i=0; i<Q.pdatalength; i++)
{
    Q.pdata[i]=0;
}

//
//
#pragma omp parallel for
for (size_t i=0; i<R.pdatalength; i++)
{
R.pdata[i]=0;
}

would then fail, even if the arrays were initialized by

vectorA4_data={12,-51,4,6,167,-68,-4,24,-41};
vectorQ4_data(9,0);
vectorR4_data(9,0);
size_t rows4 = 3, cols4 = 3;

mdspan<double, std::vector<size_t>> A4(A4_data.data(), true, {rows4, cols4});
mdspan<double, std::vector<size_t>> Q4(Q4_data.data(), true, {rows4, cols4});
mdspan<double, std::vector<size_t>> R4(R4_data.data(), true, {rows4, cols4});

and the constructors of mdspan initialize the pdatastruct like this:

template <typename T, typename Container>
mdspan<T, Container>::mdspan(T* data, const bool rowm,const Container& extents)
: pdatastruct(data,0,rowm,extents.size(),nullptr,nullptr,false,false),
pownsdata(false),
pwith_memmap(false)
{
initialize_extents(extents);
compute_strides(pextents,pstrides,rowm);
pdatastruct.pstrides = pstrides.data();
pdatastruct.pdatalength=compute_data_length(pdatastruct.pextents,pdatastruct.pstrides,pdatastruct.prank);
}
template <typename T, typename Container>
void mdspan<T, Container>::initialize_extents(const Container& extents)
{
const size_t r = extents.size();
if constexpr (StaticContainer)
{
pextents = {};
}

if constexpr (DynamicContainer<Container>)
{
    pextents.resize(r);

}

#pragma omp simd
for (size_t i = 0; i < r; ++i)
{
    pextents[i] = extents[i];
}
// Assign to datastruct
pdatastruct.pextents = pextents.data();

}

template
void compute_strides(const Container& extents, Container& strides,const bool rowmajor)
{
const size_t n = extents.size();

if constexpr (StaticContainer<Container>)
{
    strides = {}; // Default-initialize static container
}

if constexpr (DynamicContainer<Container>)
{
    strides.resize(n); // Resize dynamic container
}

if (rowmajor)
{
    // Row-major layout: last dimension has stride 1
    strides[n - 1] = 1;
    for (int i = n - 2; i >= 0; --i)
    {
        strides[i] = strides[i + 1] * extents[i + 1];
    }
}
else
{
    // Column-major layout: first dimension has stride 1
    strides[0] = 1;
    for (size_t i = 1; i < n; ++i)
    {
        strides[i] = strides[i - 1] * extents[i - 1];
    }
}

}

#pragma omp begin declare target
inline size_t compute_data_length(const size_t*__restrict extents, const size_t*__restrict strides,const size_t rank)
{
size_t offset=0;
//#pragma acc loop vector independent reduction(+:offset)
for (size_t i = 0; i < rank; ++i)
{
offset += (extents[i]-1) * strides[i];
}
return offset+1;
}
#pragma omp end declare target

So at the end of that, the arrays of the datastruct which is mapped have correctly filled members pdatalength, and prank and the pointers to the arrays pextents and pstrides and pdata are also correctly filled.

this is evicenced by the fact that without optimization, the lu decomposition works, but it is called like this:

const datastructdA=A.pdatastruct;

    datastruct<T> dL=L.pdatastruct, dU=U.pdatastruct;
    T *__restrict buffer=(T*) omp_target_alloc(sizeof(T)*2*dA.pdatalength,omp_get_default_device());
    create_in_struct(dA);
    create_out_struct(dL);
    create_out_struct(dU);

    #pragma omp target map(to:step_size) is_device_ptr(buffer)
    {
        gpu_lu_decomposition( dA,  dL, dU, (T*) buffer,step_size);
    }

the code that calls the qr decomposition is called like this:

 datastruct<T> dA=A.pdatastruct;
    datastruct<T> dQ=Q.pdatastruct;
    datastruct<T> dR=R.pdatastruct;
    T* __restrict buffer=(T*) omp_target_alloc(sizeof(T)*(dA.pextents[1]*dA.pextents[1]+2*dA.pextents[0]* dA.pextents[1]),omp_get_default_device());

    create_in_struct(dA);
    create_out_struct(dQ);
    create_out_struct(dR);

    #pragma omp target map(to:step_size) is_device_ptr(buffer)
    {
        gpu_qr_decomposition(dA,dQ,dR, (T*)buffer,step_size);
    }
    update_host(dQ);
    update_host(dR);

    exit_struct(dR);
    exit_struct(dQ);
    exit_struct(dA);

yet, the qr decomposition fails without optimization at this stage: #pragma omp parallel for

#pragma omp parallel for
for (size_t i=0; i<Q.pdatalength; i++)
{
    Q.pdata[i]=0;
}

//

which is similarly there in the lu decomposition...

and when i put on -O1, then the lu decomposition also fails...

to me this indicates a compiler problem, since the failures are depending at least somewhat at the optimization level..

But yes, I would also like to have a smaller reproducer.

Interestingly, the simple matrix multiplication always works. So it may have to do with complex structs...

[bschulz81]

attached is a bisect 1000 log. then even the simple matrix multiplication fails...

bisect1000log.txt

that is then the output:

Ordinary matrix multiplication, on gpu
"PluginInterface" error: Failure to synchronize stream (nil): Error in cuStreamSynchronize: an illegal memory access was encountered
omptarget error: Consult https://openmp.llvm.org/design/Runtimes.html for debugging options.
mdspan_acc.h:4021:1: omptarget fatal error 1: failure of target construct while offloading is mandatory

Process returned -1 (0xFFFFFFFF) execution time : 0.412 s
Press ENTER to continue.

So with bisect 1000 not even the matrix multiplication works,

i do not know where exactly it is when the lu decomposition stops..
but unfortunately, the qr decomposition does not work without any optimization set in...

[bschulz81]

hi the bisect limit where the matrix multiplication ceases to work is 857, i.e. it works with 856.

attached is the logfile for 857 where it fails at the simple matrix multiplication, and 856 where it works up to the qr decomposition

bisect-857.txt

bisect856.txt

here is the output at 856 where only the qr decomposition fails:

output856.txt

here is the output at 857 where even the matrix multiplication fails:

bisect-857.txt

However, as we see, the qr decomposition fails without any bisection.

But since the error is the same but the code actually very similar to the one which fails only at higher optimization levels,

I pretty much suspect the failure of the qr decomposition is also a compiler bug.

But yes, i would also have a more simple reproducer....

[bschulz81]

now that we have some bisection that introduces failures, what bitcode should i check?
it has created several files. should I upload all of them?

I can also make the reproducer smaller by just including the code part without the qr decomposition (i.e. which runs then until the bisection)...

Simply comment the lines 143-150 out of the acc_main.cpp. Then it does not run the qr_decomposition on the device, and will fail at the bisection 857 and run otherwise for smaller bisection values.

note that this is c++20 code with concepts. one needs the compile options:

-std=c++20 -g -O1 -Xarch_device -mllvm=-opt-bisect-limit=857 -save-temps -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -Wall

and link to rt m c stdc++ mpi , i.e one can try to compile it with:

clang++ ./main_acc.cpp -std=c++20 -g -O1 -Xarch_device -mllvm=-opt-bisect-limit=857 -save-temps -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -Wall -lrt -lm -lc -lstdc++ -lmpi

But why the qr decomposition fails, I don't know.

its probably related to this similar optimization bug or to the offloader, i guess...

[bschulz81]

that is what is doing at these lines that introduce the problem for the matrix multiplication, by the way:

BISECT: running pass (856) EarlyCSEPass on __kmpc_distribute_static_fini
BISECT: running pass (857) OpenMPOptPass on [module]
BISECT: NOT running pass (858) IPSCCPPass on [module]
BISECT: NOT running pass (859) CalledValuePropagationPass on [module]

once it runs pass at 857 it crashes...