The -fopenmp-simd enables the "#pragma omp simd" directive (only) as a
vectorization hint, but it did not work for SIMT vectorization.

This patch enables the feature for backends in which the TARGET_SIMT_VF hook is
present and indicates SIMT is available.  Only the NVPTX backend actually does
this, at present.

I've tested this with some simple testcases, and run the code using
"nvptx-none-run -L32 a.out" successfully, but I've not messed with any of the
testsuite effective target settings, yet, so there are not changes to the test
results.

The purpose is to allow me to test some SIMT features without using a full
offloading toolchain.

gcc/ChangeLog:

        * omp-expand.cc (expand_omp_simd): Use target.simt.vf to enable
        -fopenmp-simd on SIMT architectures.
---
 gcc/omp-expand.cc | 8 ++++++++
 1 file changed, 8 insertions(+)

diff --git a/gcc/omp-expand.cc b/gcc/omp-expand.cc
index 6acfd834118..8f7af80cd99 100644
--- a/gcc/omp-expand.cc
+++ b/gcc/omp-expand.cc
@@ -6588,6 +6588,14 @@ expand_omp_simd (struct omp_region *region, struct 
omp_for_data *fd)
 
   bool is_simt = omp_find_clause (gimple_omp_for_clauses (fd->for_stmt),
                                  OMP_CLAUSE__SIMT_);
+
+  /* Allow compatible targets to use SIMT with -fopenmp-simd.  */
+  if (!is_simt
+      && flag_openmp_simd
+      && targetm.simt.vf
+      && targetm.simt.vf () > 1)
+    is_simt = true;
+
   if (is_simt)
     {
       cfun->curr_properties &= ~PROP_gimple_lomp_dev;
-- 
2.51.0

Reply via email to