public inbox for gcc-bugs@sourceware.org
help / color / mirror / Atom feed
* [Bug target/99932] New: OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04
@ 2021-04-06 10:45 tschwinge at gcc dot gnu.org
2021-04-22 13:17 ` [Bug target/99932] " vries at gcc dot gnu.org
` (15 more replies)
0 siblings, 16 replies; 17+ messages in thread
From: tschwinge at gcc dot gnu.org @ 2021-04-06 10:45 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=99932
Bug ID: 99932
Summary: OpenACC/nvptx offloading execution regressions
starting with CUDA 11.2-era Nvidia Driver 460.27.04
Product: gcc
Version: 11.0
Status: UNCONFIRMED
Keywords: openacc
Severity: normal
Priority: P3
Component: target
Assignee: unassigned at gcc dot gnu.org
Reporter: tschwinge at gcc dot gnu.org
CC: vries at gcc dot gnu.org
Target Milestone: ---
Target: nvptx
We're seeing OpenACC/nvptx offloading execution regressions (including a lot of
timeouts) starting with CUDA 11.2-era Nvidia Driver 460.27.04. Confirmed with:
CUDA 11.2-era 460.27.04, 460.32.03, 460.39, 460.56, 460.67, and CUDA 11.3-era
465.19.01, across several variants of GPU hardware.
Explicitly (re-)confirmed good are older versions such as CUDA 9.1-era 390.12,
and CUDA 11.1-era 455.38, 455.45.01.
Most of these are in the 'vector_length > 32' testcases, but also a few others.
@@ -6147,7 +6147,7 @@ PASS:
libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims.c -DACC_DEVICE_T
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 (test
for warnings, line 596)
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 (test
for warnings, line 618)
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 (test
for excess errors)
[-PASS:-]{+FAIL:+}
libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2
execution test
@@ -6581,7 +6581,8 @@ PASS:
libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-1.c -DACC_DE
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-10.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 (test
for excess errors)
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-10.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0
execution test
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-10.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 (test
for excess errors)
[-PASS:-]{+WARNING: program timed out.+}
{+FAIL:+}
libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-10.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2
execution test
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-2.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 (test
for excess errors)
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-2.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0
execution test
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-2.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 output
pattern test
@@ -6599,32 +6600,32 @@ PASS:
libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-3.c -DACC_DE
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-3.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 output
pattern test
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-3.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2
scan-offload-tree-dump oaccdevlow "__attribute__\\(\\(oacc function \\(1, 1,
32\\)"
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-4.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 (test
for excess errors)
[-PASS:-]{+WARNING: program timed out.+}
{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-4.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0
execution test[-PASS:
libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-4.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 output
pattern test-]
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-4.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0
scan-offload-tree-dump oaccdevlow "__attribute__\\(\\(oacc function \\(1, 2,
128\\)"
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-4.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 (test
for excess errors)
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-4.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2
execution test
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-4.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 output
pattern test
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-4.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2
scan-offload-tree-dump oaccdevlow "__attribute__\\(\\(oacc function \\(1, 2,
128\\)"
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-5.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 (test
for excess errors)
[-PASS:-]{+WARNING: program timed out.+}
{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-5.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0
execution test[-PASS:
libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-5.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 output
pattern test-]
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-5.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0
scan-offload-tree-dump oaccdevlow "__attribute__\\(\\(oacc function \\(1, 2,
128\\)"
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-5.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 (test
for excess errors)
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-5.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2
execution test
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-5.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 output
pattern test
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-5.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2
scan-offload-tree-dump oaccdevlow "__attribute__\\(\\(oacc function \\(1, 2,
128\\)"
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-6.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 (test
for excess errors)
[-PASS:-]{+WARNING: program timed out.+}
{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-6.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0
execution test[-PASS:
libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-6.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 output
pattern test-]
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-6.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0
scan-offload-tree-dump oaccdevlow "__attribute__\\(\\(oacc function \\(1, 0,
128\\)"
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-6.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 (test
for excess errors)
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-6.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2
execution test
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-6.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 output
pattern test
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-6.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2
scan-offload-tree-dump oaccdevlow "__attribute__\\(\\(oacc function \\(1, 0,
128\\)"
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-7.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 (test
for excess errors)
[-PASS:-]{+WARNING: program timed out.+}
{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-7.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0
execution test[-PASS:
libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-7.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 output
pattern test-]
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-7.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0
scan-offload-tree-dump oaccdevlow "__attribute__\\(\\(oacc function \\(1, 0,
128\\)"
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-7.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 (test
for excess errors)
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-7.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2
execution test
@@ -6632,7 +6633,8 @@ PASS:
libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-7.c -DACC_DE
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-7.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2
scan-offload-tree-dump oaccdevlow "__attribute__\\(\\(oacc function \\(1, 0,
128\\)"
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-64-1.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 (test
for warnings, line 7)
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-64-1.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 (test
for excess errors)
[-PASS:-]{+WARNING: program timed out.+}
{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-64-1.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0
execution test
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-64-1.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 (test
for warnings, line 7)
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-64-1.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 (test
for excess errors)
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-64-1.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2
execution test
@@ -6643,7 +6645,8 @@ PASS:
libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-64-2.c -DACC_DEV
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-64-2.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2
execution test
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-64-2.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 output
pattern test
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-64-3.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 (test
for excess errors)
[-PASS:-]{+WARNING: program timed out.+}
{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-64-3.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0
execution test
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-64-3.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 (test
for excess errors)
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-64-3.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2
execution test
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vred2d-128.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 (test
for excess errors)
[-PASS:-]{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/vred2d-128.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0
execution test
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vred2d-128.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 (test
for excess errors)
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vred2d-128.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2
execution test
Similar for 'libgomp.oacc-c++', plus:
PASS: libgomp.oacc-c++/ref-1.C -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 (test for excess errors)
PASS: libgomp.oacc-c++/ref-1.C -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test
PASS: libgomp.oacc-c++/ref-1.C -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 (test for excess errors)
[-PASS:-]{+FAIL:+} libgomp.oacc-c++/ref-1.C -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 execution test
PASS: libgomp.oacc-fortran/gemm-2.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 (test for excess errors)
[-PASS:-]{+WARNING: program timed out.+}
{+FAIL:+} libgomp.oacc-fortran/gemm-2.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test
PASS: libgomp.oacc-fortran/gemm-2.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O1 (test for excess errors)
PASS: libgomp.oacc-fortran/gemm-2.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O1 execution test
PASS: libgomp.oacc-fortran/gemm-2.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 (test for excess errors)
@@ -10957,7 +10964,8 @@ PASS: libgomp.oacc-fortran/gemm-2.f90
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARE
PASS: libgomp.oacc-fortran/gemm-2.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -Os (test for excess errors)
PASS: libgomp.oacc-fortran/gemm-2.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -Os execution test
PASS: libgomp.oacc-fortran/gemm.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 (test for excess errors)
[-PASS:-]{+WARNING: program timed out.+}
{+FAIL:+} libgomp.oacc-fortran/gemm.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test
PASS: libgomp.oacc-fortran/gemm.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O1 (test for excess errors)
PASS: libgomp.oacc-fortran/gemm.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O1 execution test
PASS: libgomp.oacc-fortran/gemm.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 (test for excess errors)
@@ -12079,7 +12087,7 @@ PASS: libgomp.oacc-fortran/parallel-dims.f90
-DACC_DEVICE_TYPE_nvidia=1 -DACC_ME
PASS: libgomp.oacc-fortran/parallel-dims.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O1 (test for warnings, line 61)
PASS: libgomp.oacc-fortran/parallel-dims.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O1 (test for warnings, line 86)
PASS: libgomp.oacc-fortran/parallel-dims.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O1 (test for excess errors)
[-PASS:-]{+FAIL:+} libgomp.oacc-fortran/parallel-dims.f90
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O1
execution test
PASS: libgomp.oacc-fortran/parallel-dims.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 (test for warnings, line 61)
PASS: libgomp.oacc-fortran/parallel-dims.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 (test for warnings, line 86)
PASS: libgomp.oacc-fortran/parallel-dims.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 (test for excess errors)
@@ -12095,7 +12103,7 @@ PASS: libgomp.oacc-fortran/parallel-dims.f90
-DACC_DEVICE_TYPE_nvidia=1 -DACC_ME
PASS: libgomp.oacc-fortran/parallel-dims.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -Os (test for warnings, line 61)
PASS: libgomp.oacc-fortran/parallel-dims.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -Os (test for warnings, line 86)
PASS: libgomp.oacc-fortran/parallel-dims.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -Os (test for excess errors)
[-PASS:-]{+FAIL:+} libgomp.oacc-fortran/parallel-dims.f90
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -Os
execution test
PASS: libgomp.oacc-fortran/routine-7.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 (test for excess errors)
[-PASS:-]{+FAIL:+} libgomp.oacc-fortran/routine-7.f90
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0
execution test
PASS: libgomp.oacc-fortran/routine-7.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O1 (test for excess errors)
PASS: libgomp.oacc-fortran/routine-7.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O1 execution test
PASS: libgomp.oacc-fortran/routine-7.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 (test for excess errors)
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/99932] OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04
2021-04-06 10:45 [Bug target/99932] New: OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04 tschwinge at gcc dot gnu.org
@ 2021-04-22 13:17 ` vries at gcc dot gnu.org
2021-04-23 8:35 ` vries at gcc dot gnu.org
` (14 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2021-04-22 13:17 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=99932
--- Comment #1 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Thomas Schwinge from comment #0)
> We're seeing OpenACC/nvptx offloading execution regressions (including a lot
> of timeouts) starting with CUDA 11.2-era Nvidia Driver 460.27.04. Confirmed
> with: CUDA 11.2-era 460.27.04, 460.32.03, 460.39, 460.56, 460.67, and CUDA
> 11.3-era 465.19.01, across several variants of GPU hardware.
>
> Explicitly (re-)confirmed good are older versions such as CUDA 9.1-era
> 390.12, and CUDA 11.1-era 455.38, 455.45.01.
>
> Most of these are in the 'vector_length > 32' testcases, but also a few
> others.
>
Confirmed, I see on ubuntu 18.04.5 with dirver version 460.67:
...
FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2
execution test
...
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/99932] OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04
2021-04-06 10:45 [Bug target/99932] New: OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04 tschwinge at gcc dot gnu.org
2021-04-22 13:17 ` [Bug target/99932] " vries at gcc dot gnu.org
@ 2021-04-23 8:35 ` vries at gcc dot gnu.org
2021-04-23 12:49 ` vries at gcc dot gnu.org
` (13 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2021-04-23 8:35 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=99932
--- Comment #2 from Tom de Vries <vries at gcc dot gnu.org> ---
Minimal example:
...
$ cat libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
int
main (void)
{
int vectors_max = -1;
#pragma acc parallel \
num_gangs (1) \
num_workers (1) \
vector_length (32) \
copy (vectors_max)
{
#pragma acc loop gang reduction (max: vectors_max)
for (int i = 0; i < 2; i++)
#pragma acc loop worker reduction (max: vectors_max)
for (int j = 0; j < 2; j++)
#pragma acc loop vector reduction (max: vectors_max)
for (int k = 0; k < 32; k++)
vectors_max = k;
}
if (vectors_max != 31)
__builtin_abort ();
return 0;
}
...
Passes with GOMP_NVPTX_JIT=-O0, starts failing at GOMP_NVPTX_JIT=-O1.
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/99932] OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04
2021-04-06 10:45 [Bug target/99932] New: OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04 tschwinge at gcc dot gnu.org
2021-04-22 13:17 ` [Bug target/99932] " vries at gcc dot gnu.org
2021-04-23 8:35 ` vries at gcc dot gnu.org
@ 2021-04-23 12:49 ` vries at gcc dot gnu.org
2021-04-23 15:45 ` vries at gcc dot gnu.org
` (12 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2021-04-23 12:49 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=99932
--- Comment #3 from Tom de Vries <vries at gcc dot gnu.org> ---
Created attachment 50660
--> https://gcc.gnu.org/bugzilla/attachment.cgi?id=50660&action=edit
Cuda reproducer
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/99932] OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04
2021-04-06 10:45 [Bug target/99932] New: OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04 tschwinge at gcc dot gnu.org
` (2 preceding siblings ...)
2021-04-23 12:49 ` vries at gcc dot gnu.org
@ 2021-04-23 15:45 ` vries at gcc dot gnu.org
2021-04-24 20:27 ` vries at gcc dot gnu.org
` (11 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2021-04-23 15:45 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=99932
--- Comment #4 from Tom de Vries <vries at gcc dot gnu.org> ---
Created attachment 50662
--> https://gcc.gnu.org/bugzilla/attachment.cgi?id=50662&action=edit
Updated cuda reproducer
Slimmed down further, eliminated gang/worker reduction parts.
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/99932] OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04
2021-04-06 10:45 [Bug target/99932] New: OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04 tschwinge at gcc dot gnu.org
` (3 preceding siblings ...)
2021-04-23 15:45 ` vries at gcc dot gnu.org
@ 2021-04-24 20:27 ` vries at gcc dot gnu.org
2021-04-27 6:58 ` vries at gcc dot gnu.org
` (10 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2021-04-24 20:27 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=99932
--- Comment #5 from Tom de Vries <vries at gcc dot gnu.org> ---
FIled https://developer.nvidia.com/nvidia_bug/3299227
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/99932] OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04
2021-04-06 10:45 [Bug target/99932] New: OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04 tschwinge at gcc dot gnu.org
` (4 preceding siblings ...)
2021-04-24 20:27 ` vries at gcc dot gnu.org
@ 2021-04-27 6:58 ` vries at gcc dot gnu.org
2021-12-10 6:46 ` vries at gcc dot gnu.org
` (9 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2021-04-27 6:58 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=99932
--- Comment #6 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Tom de Vries from comment #5)
> FIled https://developer.nvidia.com/nvidia_bug/3299227
Nvidia reported it will be fixed in the next major cuda release. I've asked
about driver fixes.
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/99932] OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04
2021-04-06 10:45 [Bug target/99932] New: OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04 tschwinge at gcc dot gnu.org
` (5 preceding siblings ...)
2021-04-27 6:58 ` vries at gcc dot gnu.org
@ 2021-12-10 6:46 ` vries at gcc dot gnu.org
2022-01-24 9:59 ` vries at gcc dot gnu.org
` (8 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2021-12-10 6:46 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=99932
--- Comment #7 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Tom de Vries from comment #6)
> (In reply to Tom de Vries from comment #5)
> > FIled https://developer.nvidia.com/nvidia_bug/3299227
>
> Nvidia reported it will be fixed in the next major cuda release. I've asked
> about driver fixes.
I've tested 470.86, and can confirm that the failure as reported to nvidia has
disappeared.
However, the example from comment 2 still passes with GOMP_NVPTX_JIT=-O0, and
starts failing at GOMP_NVPTX_JIT=-O1.
So it looks like we'll need to file another bug-report at nvidia.
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/99932] OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04
2021-04-06 10:45 [Bug target/99932] New: OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04 tschwinge at gcc dot gnu.org
` (6 preceding siblings ...)
2021-12-10 6:46 ` vries at gcc dot gnu.org
@ 2022-01-24 9:59 ` vries at gcc dot gnu.org
2022-01-24 10:04 ` vries at gcc dot gnu.org
` (7 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2022-01-24 9:59 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=99932
--- Comment #8 from Tom de Vries <vries at gcc dot gnu.org> ---
New minimal oacc example:
...
int
main (void)
{
int vectors_max = -1;
#pragma acc parallel \
num_gangs (1) num_workers (1) \
copy (vectors_max)
{
for (int i = 0; i < 2; i++)
for (int j = 0; j < 2; j++)
#pragma acc loop vector reduction (max: vectors_max)
for (int k = 0; k < 32; k++)
vectors_max = k;
}
if (vectors_max != 31)
__builtin_abort ();
return 0;
}
...
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/99932] OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04
2021-04-06 10:45 [Bug target/99932] New: OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04 tschwinge at gcc dot gnu.org
` (7 preceding siblings ...)
2022-01-24 9:59 ` vries at gcc dot gnu.org
@ 2022-01-24 10:04 ` vries at gcc dot gnu.org
2022-01-24 15:46 ` vries at gcc dot gnu.org
` (6 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2022-01-24 10:04 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=99932
--- Comment #9 from Tom de Vries <vries at gcc dot gnu.org> ---
Created attachment 52273
--> https://gcc.gnu.org/bugzilla/attachment.cgi?id=52273&action=edit
New cuda reproducer
$ ./do.sh
DRIVER SASS, ptxas=-O0:
+ /home/vries/cuda/11.4.3/bin/nvcc vector-max.cu -Wno-deprecated-gpu-targets
-arch=compute_35 -Xptxas -O0
+ ./a.out
a[0]: 31
DRIVER SASS, ptxas=-O1:
+ /home/vries/cuda/11.4.3/bin/nvcc vector-max.cu -Wno-deprecated-gpu-targets
-arch=compute_35 -Xptxas -O1
+ ./a.out
a[0]: 0
./do.sh: line 34: 27353 Aborted (core dumped) ./a.out
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/99932] OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04
2021-04-06 10:45 [Bug target/99932] New: OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04 tschwinge at gcc dot gnu.org
` (8 preceding siblings ...)
2022-01-24 10:04 ` vries at gcc dot gnu.org
@ 2022-01-24 15:46 ` vries at gcc dot gnu.org
2022-01-25 10:42 ` vries at gcc dot gnu.org
` (5 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2022-01-24 15:46 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=99932
--- Comment #10 from Tom de Vries <vries at gcc dot gnu.org> ---
[ FTR, T400, driver 470.94 ]
Interestingly, changing the default ptx version to 6.3 makes the minimal
test-case pass, as well as the full parallel-dims.c
The only code changes are shfl -> shfl.sync and vote -> vote.sync.
Both of these require sm_30, so from that perspective we won't leave any
architectures behind.
OTOH, this may leave behind:
- some older drivers
- some older CUDAs (if ptxas is used for ptx verification in the
nvptx-none-as).
Rerunning the entire testsuite though shows that the non-32-vector-length
test-cases are still failing.
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/99932] OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04
2021-04-06 10:45 [Bug target/99932] New: OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04 tschwinge at gcc dot gnu.org
` (9 preceding siblings ...)
2022-01-24 15:46 ` vries at gcc dot gnu.org
@ 2022-01-25 10:42 ` vries at gcc dot gnu.org
2022-01-25 11:02 ` vries at gcc dot gnu.org
` (4 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2022-01-25 10:42 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=99932
--- Comment #11 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Tom de Vries from comment #10)
> Rerunning the entire testsuite though shows that the non-32-vector-length
> test-cases are still failing.
Minimal example:
...
int
main (void)
{
#pragma acc parallel num_workers (2) vector_length (64)
{
#pragma acc loop worker
for (unsigned int i = 0; i < 2; i++)
#pragma acc loop vector
for (unsigned int j = 0; j < 64; j++)
;
}
return 0;
}
...
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/99932] OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04
2021-04-06 10:45 [Bug target/99932] New: OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04 tschwinge at gcc dot gnu.org
` (10 preceding siblings ...)
2022-01-25 10:42 ` vries at gcc dot gnu.org
@ 2022-01-25 11:02 ` vries at gcc dot gnu.org
2022-01-25 15:00 ` vries at gcc dot gnu.org
` (3 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2022-01-25 11:02 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=99932
--- Comment #12 from Tom de Vries <vries at gcc dot gnu.org> ---
Created attachment 52285
--> https://gcc.gnu.org/bugzilla/attachment.cgi?id=52285&action=edit
Cuda reproducer non-32 vector length
[ On T400, driver version 470.94 ]
NVCC SASS:
...
$ ./do.sh
NVCC SASS, ptxas=-O0:
+ /home/vries/cuda/11.5.1/bin/nvcc vector-length-64.cu -arch=compute_75
-code=sm_75 -Xptxas -O0
+ ./a.out
NVCC SASS, ptxas=-O1:
+ /home/vries/cuda/11.5.1/bin/nvcc vector-length-64.cu -arch=compute_75
-code=sm_75 -Xptxas -O1
+ ./a.out
<HANGS>
...
Driver SASS:
...
$ ./do.sh
DRIVER SASS, ptxas=-O0:
+ /home/vries/cuda/11.4.3/bin/nvcc vector-length-64.cu -arch=compute_75 -Xptxas
-O0
+ ./a.out
DRIVER SASS, ptxas=-O1:
+ /home/vries/cuda/11.4.3/bin/nvcc vector-length-64.cu -arch=compute_75 -Xptxas
-O1
+ ./a.out
<HANGS>
...
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/99932] OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04
2021-04-06 10:45 [Bug target/99932] New: OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04 tschwinge at gcc dot gnu.org
` (11 preceding siblings ...)
2022-01-25 11:02 ` vries at gcc dot gnu.org
@ 2022-01-25 15:00 ` vries at gcc dot gnu.org
2022-01-26 12:19 ` vries at gcc dot gnu.org
` (2 subsequent siblings)
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2022-01-25 15:00 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=99932
--- Comment #13 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Tom de Vries from comment #10)
> [ FTR, T400, driver 470.94 ]
>
> Interestingly, changing the default ptx version to 6.3 makes the minimal
> test-case pass, as well as the full parallel-dims.c
>
> The only code changes are shfl -> shfl.sync and vote -> vote.sync.
>
It seems another change is required.
Starting with 6.0, bar.sync maps onto barrier.sync.aligned, where the aligned
means that "all threads in CTA will execute the same barrier instruction. In
conditionally executed code, an aligned barrier instruction should only be used
if it is known that all threads in CTA evaluate the condition identically,
otherwise behavior is undefined."
It's not fully clear what is meant with "the same barrier instruction" or
"condition", but in the case of vector_length > 32, we use:
...
bar.sync %r67,64;
...
where %r67 is a barrier number, 1 for worker 0 and 2 for worker 1 in case of 2
workers. It may well be that it's invalid to use bar.sync for this, and we
should use barrier.sync instead.
But then there's an isa note:
...
Note: For .target sm_6x or below,
1. barrier instruction without .aligned modifier is equivalent to .aligned
variant and has the same restrictions as of .aligned variant.
...
which seems to imply that we get back barrier.sync.aligned behaviour for sm_6x
and earlier, which would again break vector_length > 32.
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/99932] OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04
2021-04-06 10:45 [Bug target/99932] New: OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04 tschwinge at gcc dot gnu.org
` (12 preceding siblings ...)
2022-01-25 15:00 ` vries at gcc dot gnu.org
@ 2022-01-26 12:19 ` vries at gcc dot gnu.org
2022-01-26 12:30 ` vries at gcc dot gnu.org
2022-02-02 12:31 ` vries at gcc dot gnu.org
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2022-01-26 12:19 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=99932
--- Comment #14 from Tom de Vries <vries at gcc dot gnu.org> ---
An observation when playing around with vector-length-128-4.c: there are two
ways in which I can make the example pass.
1. add barrier.sync.aligned 0 or membar.cta after first broad-cast receive
2. unroll loop in first broad-cast send.
At first glance, it doesn't look entirely trivial though to implement either.
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/99932] OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04
2021-04-06 10:45 [Bug target/99932] New: OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04 tschwinge at gcc dot gnu.org
` (13 preceding siblings ...)
2022-01-26 12:19 ` vries at gcc dot gnu.org
@ 2022-01-26 12:30 ` vries at gcc dot gnu.org
2022-02-02 12:31 ` vries at gcc dot gnu.org
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2022-01-26 12:30 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=99932
--- Comment #15 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Tom de Vries from comment #14)
> An observation when playing around with vector-length-128-4.c:
Another observation:
...
$L11:
ld.u64 %r108,[%r109];
st.u64 [%r112],%r108;
setp.lt.u32 %r111,%r110,3;
add.u32 %r110,%r110,1;
add.u64 %r109,%r109,8;
add.u64 %r112,%r112,8;
@ %r111 bra.uni $L11;
...
The bra.uni in the broadcast loop is incorrect, it's used in a vector-neutered
block.
^ permalink raw reply [flat|nested] 17+ messages in thread* [Bug target/99932] OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04
2021-04-06 10:45 [Bug target/99932] New: OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04 tschwinge at gcc dot gnu.org
` (14 preceding siblings ...)
2022-01-26 12:30 ` vries at gcc dot gnu.org
@ 2022-02-02 12:31 ` vries at gcc dot gnu.org
15 siblings, 0 replies; 17+ messages in thread
From: vries at gcc dot gnu.org @ 2022-02-02 12:31 UTC (permalink / raw)
To: gcc-bugs
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=99932
Tom de Vries <vries at gcc dot gnu.org> changed:
What |Removed |Added
----------------------------------------------------------------------------
Keywords| |testsuite-fail
Resolution|--- |FIXED
Target Milestone|--- |12.0
Status|UNCONFIRMED |RESOLVED
--- Comment #16 from Tom de Vries <vries at gcc dot gnu.org> ---
All these fails should be fixed on current trunk.
^ permalink raw reply [flat|nested] 17+ messages in thread
end of thread, other threads:[~2022-02-02 12:31 UTC | newest]
Thread overview: 17+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-04-06 10:45 [Bug target/99932] New: OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04 tschwinge at gcc dot gnu.org
2021-04-22 13:17 ` [Bug target/99932] " vries at gcc dot gnu.org
2021-04-23 8:35 ` vries at gcc dot gnu.org
2021-04-23 12:49 ` vries at gcc dot gnu.org
2021-04-23 15:45 ` vries at gcc dot gnu.org
2021-04-24 20:27 ` vries at gcc dot gnu.org
2021-04-27 6:58 ` vries at gcc dot gnu.org
2021-12-10 6:46 ` vries at gcc dot gnu.org
2022-01-24 9:59 ` vries at gcc dot gnu.org
2022-01-24 10:04 ` vries at gcc dot gnu.org
2022-01-24 15:46 ` vries at gcc dot gnu.org
2022-01-25 10:42 ` vries at gcc dot gnu.org
2022-01-25 11:02 ` vries at gcc dot gnu.org
2022-01-25 15:00 ` vries at gcc dot gnu.org
2022-01-26 12:19 ` vries at gcc dot gnu.org
2022-01-26 12:30 ` vries at gcc dot gnu.org
2022-02-02 12:31 ` vries at gcc dot gnu.org
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).