This PR fixes the too-much-OpenACC-warning issue, mentioned in PR #6639.

This PR also re-enables the OpenACC CI test.
kokkos · Dec 1, 2023 · 6856209 · 6856209
1 parent 9041bda
commit 6856209
Show file tree

Hide file tree

Showing 4 changed files with 56 additions and 44 deletions.
diff --git a/.jenkins b/.jenkins
@@ -32,32 +32,32 @@ pipeline {
         }
         stage('Build') {
             parallel {
-//                stage('OPENACC-NVHPC-CUDA-12.2') {
-//                    agent {
-//                        dockerfile {
-//                            filename 'Dockerfile.nvhpc'
-//                            dir 'scripts/docker'
-//                            label 'nvidia-docker && volta && large_images'
-//                            args '--env NVIDIA_VISIBLE_DEVICES=$NVIDIA_VISIBLE_DEVICES'
-//                        }
-//                    }
-//                    environment {
-//                        NVHPC_CUDA_HOME = '/opt/nvidia/hpc_sdk/Linux_x86_64/23.7/cuda/12.2'
-//                    }
-//                    steps {
-//                        sh '''rm -rf build && mkdir -p build && cd build && \
-//                              /opt/cmake/bin/cmake \
-//                                -DCMAKE_CXX_COMPILER=nvc++ \
-//                                -DCMAKE_CXX_STANDARD=17 \
-//                                -DKokkos_ARCH_NATIVE=ON \
-//                                -DKokkos_ENABLE_COMPILER_WARNINGS=ON \
-//                                -DKokkos_ENABLE_TESTS=ON \
-//                                -DKokkos_ENABLE_OPENACC=ON \
-//                                -DKokkos_ARCH_VOLTA70=ON \
-//                              .. && \
-//                              make -j8 && ctest --verbose'''
-//                    }
-//                }
+                stage('OPENACC-NVHPC-CUDA-12.2') {
+                    agent {
+                        dockerfile {
+                            filename 'Dockerfile.nvhpc'
+                            dir 'scripts/docker'
+                            label 'nvidia-docker && volta && large_images'
+                            args '--env NVIDIA_VISIBLE_DEVICES=$NVIDIA_VISIBLE_DEVICES'
+                        }
+                    }
+                    environment {
+                        NVHPC_CUDA_HOME = '/opt/nvidia/hpc_sdk/Linux_x86_64/23.7/cuda/12.2'
+                    }
+                    steps {
+                        sh '''rm -rf build && mkdir -p build && cd build && \
+                              /opt/cmake/bin/cmake \
+                                -DCMAKE_CXX_COMPILER=nvc++ \
+                                -DCMAKE_CXX_STANDARD=17 \
+                                -DKokkos_ARCH_NATIVE=ON \
+                                -DKokkos_ENABLE_COMPILER_WARNINGS=ON \
+                                -DKokkos_ENABLE_TESTS=ON \
+                                -DKokkos_ENABLE_OPENACC=ON \
+                                -DKokkos_ARCH_VOLTA70=ON \
+                              .. && \
+                              make -j8 && ctest --verbose'''
+                    }
+                }
                 stage('CUDA-12.2-NVHPC') {
                     agent {
                         dockerfile {

diff --git a/tpls/desul/include/desul/atomics/Compare_Exchange_OpenACC.hpp b/tpls/desul/include/desul/atomics/Compare_Exchange_OpenACC.hpp
@@ -33,9 +33,11 @@ T device_atomic_exchange(T* dest, T value, MemoryOrder, MemoryScope /*scope*/) {
     return return_val;
   } else {
     // FIXME_OPENACC
-    printf(
-        "DESUL error in device_atomic_exchange(): Not supported atomic operation in "
-        "the OpenACC backend\n");
+    if (acc_on_device(acc_device_not_host)) {
+      printf(
+          "DESUL error in device_atomic_exchange(): Not supported atomic operation in "
+          "the OpenACC backend\n");
+    }
     //  Acquire a lock for the address
     // while (!lock_address_openacc((void*)dest, scope)) {
     // }
@@ -73,9 +75,11 @@ T device_atomic_compare_exchange(
     return atomicCAS(dest, compare, value);
   } else {
     // FIXME_OPENACC
-    printf(
-        "DESUL error in device_atomic_compare_exchange(): Not supported atomic "
-        "operation in the OpenACC backend\n");
+    if (acc_on_device(acc_device_not_host)) {
+      printf(
+          "DESUL error in device_atomic_compare_exchange(): Not supported atomic "
+          "operation in the OpenACC backend\n");
+    }
     T current_val = *dest;
     // Acquire a lock for the address
     // while (!lock_address_openacc((void*)dest, scope)) {

diff --git a/tpls/desul/include/desul/atomics/Fetch_Op_OpenACC.hpp b/tpls/desul/include/desul/atomics/Fetch_Op_OpenACC.hpp
@@ -390,9 +390,11 @@ std::enable_if_t<is_openacc_arithmetic_type_v<T>, void> device_atomic_store(
 template <class T>
 std::enable_if_t<is_openacc_arithmetic_type_v<T>, void> device_atomic_store(
     T* const ptr, const T val, MemoryOrderRelease, MemoryScopeDevice) {
-  printf(
-      "DESUL error in device_atomic_store(MemoryOrderRelease): Not supported atomic "
-      "operation in the OpenACC backend\n");
+  if (acc_on_device(acc_device_not_host)) {
+    printf(
+        "DESUL error in device_atomic_store(MemoryOrderRelease): Not supported atomic "
+        "operation in the OpenACC backend\n");
+  }
 #pragma acc atomic write
   *ptr = val;
 }
@@ -411,9 +413,11 @@ std::enable_if_t<is_openacc_arithmetic_type_v<T>, T> device_atomic_load(
 template <class T>
 std::enable_if_t<is_openacc_arithmetic_type_v<T>, T> device_atomic_load(
     const T* const ptr, MemoryOrderAcquire, MemoryScopeDevice) {
-  printf(
-      "DESUL error in device_atomic_load(MemoryOrderAcquire): Not supported atomic "
-      "operation in the OpenACC backend\n");
+  if (acc_on_device(acc_device_not_host)) {
+    printf(
+        "DESUL error in device_atomic_load(MemoryOrderAcquire): Not supported atomic "
+        "operation in the OpenACC backend\n");
+  }
   T retval;
 #pragma acc atomic read
   retval = *ptr;

diff --git a/tpls/desul/include/desul/atomics/Lock_Based_Fetch_Op_OpenACC.hpp b/tpls/desul/include/desul/atomics/Lock_Based_Fetch_Op_OpenACC.hpp
@@ -29,9 +29,11 @@ inline T device_atomic_fetch_oper(const Oper& op,
                                   dont_deduce_this_parameter_t<const T> val,
                                   MemoryOrder /*order*/,
                                   MemoryScope scope) {
-  printf(
-      "DESUL error in device_atomic_fetch_oper(): Not supported atomic operation in "
-      "the OpenACC backend\n");
+  if (acc_on_device(acc_device_not_host)) {
+    printf(
+        "DESUL error in device_atomic_fetch_oper(): Not supported atomic operation in "
+        "the OpenACC backend\n");
+  }
   // Acquire a lock for the address
   while (!lock_address((void*)dest, scope)) {
   }
@@ -56,9 +58,11 @@ inline T device_atomic_oper_fetch(const Oper& op,
                                   dont_deduce_this_parameter_t<const T> val,
                                   MemoryOrder /*order*/,
                                   MemoryScope scope) {
-  printf(
-      "DESUL error in device_atomic_oper_fetch(): Not supported atomic operation in "
-      "the OpenACC backend\n");
+  if (acc_on_device(acc_device_not_host)) {
+    printf(
+        "DESUL error in device_atomic_oper_fetch(): Not supported atomic operation in "
+        "the OpenACC backend\n");
+  }
   // Acquire a lock for the address
   while (!lock_address((void*)dest, scope)) {
   }