https://gcc.gnu.org/bugzilla/show_bug.cgi?id=122281

Tobias Burnus <burnus at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
                 CC|                            |burnus at gcc dot gnu.org
           Keywords|                            |wrong-code

--- Comment #3 from Tobias Burnus <burnus at gcc dot gnu.org> ---
(In reply to Benjamin Schulz from comment #1)
> Created attachment 62558 [details]
> archiv.tar.gz
> 
> above, i postet a precompiled .ii file. Here is the full source code.
> 
> I dont know why this suddenly happens...

The file seems identical to PR 122280's first attachment, but this time the
file to test is   sparsetests.cpp  

If I compile it as:

mpicxx -std=c++23 -fopenmp -I . sparsetests.cpp -g -O3
-foffload-options=nvptx-none=-march=sm_80

using the git version of g++ and CUDA 13 + Nvidia RTX A1000 6GB Laptop.

I get the following segfault (host code):

Thread 20 "a.out" received signal SIGSEGV, Segmentation fault.
[Switching to Thread 0x7fffed2fa6c0 (LWP 3041701)]
0x0000000000406c18 in _ZNK9DataBlockIdE8sparsityEv._omp_fn.2(void) () at
/dev/shm/z/datablock.h:703
703                     #pragma omp atomic update
(gdb) list
698             #pragma omp parallel for simd shared(count)
699             for(size_t i=0; i<dpdatalength; i++)
700             {
701                 if(dpdata[i]==0)
702                 {
703                     #pragma omp atomic update
704                     count++;
705                 }

with:

(gdb) p i
$1 = 61
(gdb) p dpdatalength
$2 = 64
(gdb) p dpdata[i]
$3 = 0
(gdb) p count
$4 = 0

* * *

Compiling with -O0 solves this - but it re-appears with -O1.

That's in line with valgrind:

==3042608== Thread 9:
==3042608== Invalid read of size 8
==3042608==    at 0x4025C1: DataBlock<double>::sparsity() const [clone
._omp_fn.2] (datablock.h:703)

__attribute__((omp declare target))
void DataBlock<double>::sparsity (struct .omp_data_s.156 & restrict
.omp_data_i)
{
...
  size_t * _29(D);
...
  <bb 8> [local count: 620085900]:
  [datablock.h:703:25 discrim 3] __atomic_fetch_add_8 (_29(D), 1, 0);

* * *

Compare this to -O0:

  <bb 10> :
  [datablock.h:703:25] _32 = .omp_data_i_8(D)->count;
  [datablock.h:703:25 discrim 1] __atomic_fetch_add_8 (_32, 1, 0);

* * *

I wonder what removes the '= .omp_data_i_8(D)->count;' assignment with -O1 ?!?

Reply via email to