LLNL · adrienbernede · Aug 30, 2024 · Aug 30, 2024 · Sep 4, 2024 · Sep 18, 2024
diff --git a/.gitlab/jobs/corona.yml b/.gitlab/jobs/corona.yml
@@ -27,6 +27,9 @@
 # ${PROJECT_<MACHINE>_DEPS} in the extra jobs. There is no reason not to fully
 # describe the spec here.
 
-# With GitLab CI, included files cannot be empty.
-variables:
-  INCLUDED_FILE_CANNOT_BE_EMPTY: "True"
+clang_19_0_0_sycl_gcc_10_3_1_rocmcc_5_7_1_hip:
+  variables:
+    SPEC: " ~shared +sycl ~openmp %clang@=19.0.0 cxxflags==\"-w -fsycl -fsycl-unnamed-lambda -fsycl-targets=amdgcn-amd-amdhsa -Xsycl-target-backend --offload-arch=gfx906\" ^blt@develop"
+    MODULE_LIST: "rocm/5.7.1"
+  extends: .job_on_corona
+  allow_failure: true
diff --git a/.gitlab/jobs/lassen.yml b/.gitlab/jobs/lassen.yml
@@ -18,14 +18,7 @@
 # We keep ${PROJECT_<MACHINE>_VARIANTS} and ${PROJECT_<MACHINE>_DEPS} So that
 # the comparison with the original job is easier.
 
-# Overriding shared spec: Longer allocation + extra flags
-xl_2022_08_19_gcc_8_3_1_cuda_11_2_0:
-  variables:
-    SPEC: "${PROJECT_LASSEN_VARIANTS} +cuda cxxflags==\"-qthreaded -std=c++14 -O3 -qstrict -qxlcompatmacros -qlanglvl=extended0x -qalias=noansi -qhot -qpic -qsmp=omp -qsuppress=1500-029 -qsuppress=1500-036\" %xl@=16.1.1.12.gcc.8.3.1 ^[email protected]+allow-unsupported-compilers ${PROJECT_LASSEN_DEPS}"
-    MODULE_LIST: "cuda/11.2.0"
-    LASSEN_JOB_ALLOC: "1 -W 60 -q pci"
-  extends: .job_on_lassen
-
+# No jobs overridden
 
 ############
 # Extra jobs
@@ -36,7 +29,7 @@ xl_2022_08_19_gcc_8_3_1_cuda_11_2_0:
 
 gcc_8_3_1:
   variables:
-    SPEC: " ~shared +openmp %gcc@=8.3.1 ${PROJECT_LASSEN_DEPS}"
+    SPEC: " ~shared +openmp %gcc@=8.3.1 ^blt@develop"
   extends: .job_on_lassen
 
 gcc_8_3_1_cuda_11_5_0_ats_disabled:
@@ -69,8 +62,23 @@ clang_13_0_1_libcpp:
 #    LSAN_OPTIONS: "suppressions=${CI_PROJECT_DIR}/tpl/RAJA/suppressions.asan"
 #  extends: .job_on_lassen
 
-# Activated in RAJA, but we don't use desul atomics here
-#gcc_8_3_1_cuda_10_1_168_desul_atomics:
-#  variables:
-#    SPEC: "+openmp +cuda +desul %gcc@=8.3.1 cuda_arch=70 cuda_arch=70 ^[email protected]+allow-unsupported-compilers ${PROJECT_LASSEN_DEPS}"
-#  extends: .job_on_lassen
+clang_16_0_6_ibm_omptarget:
+  variables:
+    SPEC: " ~shared +openmp +omptarget %clang@=16.0.6.ibm.gcc.8.3.1 ^blt@develop"
+    ON_LASSEN: "OFF"
+  extends: .job_on_lassen
+
+xl_2022_08_19_gcc_8_3_1_cuda_11_2_0:
+  variables:
+    SPEC: " ~shared +openmp cuda_arch=70 +cuda cxxflags==\"-qthreaded -std=c++14 -O3 -qstrict -qxlcompatmacros -qlanglvl=extended0x -qalias=noansi -qhot -qpic -qsmp=omp -qsuppress=1500-029 -qsuppress=1500-036\" %xl@=16.1.1.12.gcc.8.3.1 ^[email protected]+allow-unsupported-compilers ^blt@develop"
+    MODULE_LIST: "cuda/11.2.0"
+    LASSEN_JOB_ALLOC: "1 -W 60 -q pci"
+  extends: .job_on_lassen
+
+xl_2023_06_28_gcc_11_2_1_cuda_11_8_0:
+  variables:
+    SPEC: " ~shared +openmp cuda_arch=70 +cuda cxxflags==\"-qthreaded -std=c++14 -O3 -qstrict -qxlcompatmacros -qlanglvl=extended0x -qalias=noansi -qhot -qpic -qsmp=omp -qsuppress=1500-029 -qsuppress=1500-036\" %xl@=16.1.1.14.cuda.11.8.0.gcc.11.2.1 ^[email protected]+allow-unsupported-compilers ^blt@develop"
+    MODULE_LIST: "cuda/11.8.0"
+    LASSEN_JOB_ALLOC: "1 -W 60 -q pci"
+  extends: .job_on_lassen
+
diff --git a/.gitlab/jobs/poodle.yml b/.gitlab/jobs/poodle.yml
@@ -18,27 +18,25 @@
 # We keep ${PROJECT_<MACHINE>_VARIANTS} and ${PROJECT_<MACHINE>_DEPS} So that
 # the comparison with the original job is easier.
 
+# allow failure
+intel_2023_2_1:
+  variables:
+    SPEC: "${PROJECT_POODLE_VARIANTS} %intel@=2023.2.1 ${PROJECT_POODLE_DEPS}"
+  extends: .job_on_poodle
+  allow_failure: true
+
+# omptask variant
 clang_14_0_6:
   variables:
     SPEC: "${PROJECT_POODLE_VARIANTS} +omptask %clang@=14.0.6 ${PROJECT_POODLE_DEPS}"
   extends: .job_on_poodle
 
+# omptask variant
 gcc_10_3_1:
   variables:
     SPEC: "${PROJECT_POODLE_VARIANTS} +omptask %gcc@=10.3.1 ${PROJECT_POODLE_DEPS}"
   extends: .job_on_poodle
 
-intel_19_1_2_gcc_10_3_1:
-  variables:
-    SPEC: "${PROJECT_POODLE_VARIANTS} %intel@=19.1.2.gcc.10.3.1 ${PROJECT_POODLE_DEPS}"
-  extends: .job_on_poodle
-
-intel_2022_1_0:
-  variables:
-    SPEC: "${PROJECT_POODLE_VARIANTS} %intel@=2022.1.0 ${PROJECT_POODLE_DEPS}"
-  allow_failure: true
-  extends: .job_on_poodle
-
 ############
 # Extra jobs
 ############
@@ -49,5 +47,4 @@ intel_2022_1_0:
 intel_2022_1_0_mpi:
   variables:
     SPEC: "~shared +openmp +mpi %intel@=2022.1.0 ^mvapich2 ^blt@develop"
-  allow_failure: true
   extends: .job_on_poodle
diff --git a/.gitlab/jobs/ruby.yml b/.gitlab/jobs/ruby.yml
@@ -18,28 +18,26 @@
 # We keep ${PROJECT_<MACHINE>_VARIANTS} and ${PROJECT_<MACHINE>_DEPS} So that
 # the comparison with the original job is easier.
 
+# allow failure
+intel_2023_2_1:
+  variables:
+    SPEC: "${PROJECT_RUBY_VARIANTS} %intel@=2023.2.1 ${PROJECT_RUBY_DEPS}"
+  extends: .job_on_ruby
+  allow_failure: true
+
+# omptask variant
 clang_14_0_6:
   variables:
     SPEC: "${PROJECT_RUBY_VARIANTS} +omptask %clang@=14.0.6 ${PROJECT_RUBY_DEPS}"
   extends: .job_on_ruby
 
+# omptask variant
 gcc_10_3_1:
   variables:
     SPEC: "${PROJECT_RUBY_VARIANTS} +omptask %gcc@=10.3.1 ${PROJECT_RUBY_DEPS}"
     RUBY_BUILD_AND_TEST_JOB_ALLOC: "--time=60 --nodes=1"
   extends: .job_on_ruby
 
-intel_19_1_2_gcc_10_3_1:
-  variables:
-    SPEC: "${PROJECT_RUBY_VARIANTS} %intel@=19.1.2.gcc.10.3.1 ${PROJECT_RUBY_DEPS}"
-    RUBY_BUILD_AND_TEST_JOB_ALLOC: "--time=40 --nodes=1"
-  extends: .job_on_ruby
-
-intel_2022_1_0:
-  variables:
-    SPEC: "${PROJECT_RUBY_VARIANTS} %intel@=2022.1.0 ${PROJECT_RUBY_DEPS}"
-  extends: .job_on_ruby
-
 ############
 # Extra jobs
 ############

diff --git a/.gitlab/jobs/tioga.yml b/.gitlab/jobs/tioga.yml
@@ -27,13 +27,13 @@
 # ${PROJECT_<MACHINE>_DEPS} in the extra jobs. There is no reason not to fully
 # describe the spec here.
 
-rocmcc_6_1_1_hip_openmp:
+rocmcc_6_2_0_hip_openmp:
   variables:
-    SPEC: "~shared +rocm +openmp amdgpu_target=gfx90a %rocmcc@=6.1.1 ^hip@6.1.1 ^blt@develop"
+    SPEC: "~shared +rocm +openmp amdgpu_target=gfx90a %rocmcc@=6.2.0 ^hip@6.2.0 ^blt@develop"
   extends: .job_on_tioga
 
-rocmcc_6_1_1_hip_openmp_mpi:
+rocmcc_6_2_0_hip_openmp_mpi:
   variables:
-    SPEC: "~shared +rocm +openmp +mpi amdgpu_target=gfx90a %rocmcc@=6.1.1 ^hip@6.1.1 ^blt@develop"
+    SPEC: "~shared +rocm +openmp +mpi amdgpu_target=gfx90a %rocmcc@=6.2.0 ^hip@6.2.0 ^blt@develop"
   extends: .job_on_tioga
   allow_failure: true
diff --git a/.uberenv_config.json b/.uberenv_config.json
@@ -1,10 +1,10 @@
 {
-"package_name" : "raja_perf",
+"package_name" : "raja-perf",
 "package_version" : "develop",
 "package_final_phase" : "initconfig",
 "package_source_dir" : "../..",
 "spack_url": "https://github.com/spack/spack.git",
-"spack_branch": "develop-2024-05-26",
+"spack_branch": "develop-2024-10-06",
 "spack_activate" : {},
 "spack_configs_path": "tpl/RAJA/scripts/radiuss-spack-configs",
 "spack_packages_path": "tpl/RAJA/scripts/radiuss-spack-configs/packages",

diff --git a/scripts/gitlab/build_and_test.sh b/scripts/gitlab/build_and_test.sh
@@ -28,6 +28,7 @@ job_unique_id=${CI_JOB_ID:-""}
 use_dev_shm=${USE_DEV_SHM:-true}
 spack_debug=${SPACK_DEBUG:-false}
 debug_mode=${DEBUG_MODE:-false}
+push_to_registry=${PUSH_TO_REGISTRY:-true}
 
 raja_version=${UPDATE_RAJA:-""}
 sys_type=${SYS_TYPE:-""}
@@ -59,6 +60,7 @@ then
     echo "~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~"
     use_dev_shm=false
     spack_debug=true
+    push_to_registry=false
 fi
 
 if [[ -n ${module_list} ]]
@@ -149,7 +151,7 @@ then
     timed_message "Spack build of dependencies"
     ${uberenv_cmd} --skip-setup-and-env --spec="${spec}" ${prefix_opt} ${upstream_opt}
 
-    if [[ -n ${ci_registry_token} && ${debug_mode} == false ]]
+    if [[ -n ${ci_registry_token} && ${push_to_registry} == true ]]
     then
         timed_message "Push dependencies to buildcache"
         ${spack_cmd} -D ${spack_env_path} buildcache push --only dependencies gitlab_ci

diff --git a/src/algorithm/ATOMIC.hpp b/src/algorithm/ATOMIC.hpp
@@ -74,6 +74,7 @@ class ATOMIC : public KernelBase
   void setOpenMPTuningDefinitions(VariantID vid);
   void setCudaTuningDefinitions(VariantID vid);
   void setHipTuningDefinitions(VariantID vid);
+  void setOpenMPTargetTuningDefinitions(VariantID vid);
 
   template < size_t replication >
   void runSeqVariantReplicate(VariantID vid);

diff --git a/src/algorithm/REDUCE_SUM-Cuda.cpp b/src/algorithm/REDUCE_SUM-Cuda.cpp
@@ -239,7 +239,8 @@ void REDUCE_SUM::runCudaVariantRAJANewReduce(VariantID vid)
       RAJA::forall<exec_policy>( res,
         RAJA::RangeSegment(ibegin, iend),
         RAJA::expt::Reduce<RAJA::operators::plus>(&tsum),
-        [=] __device__ (Index_type i, Real_type& sum) {
+        [=] __device__ (Index_type i,
+          RAJA::expt::ValOp<Real_type, RAJA::operators::plus>& sum) {
           REDUCE_SUM_BODY;
         }
       );

diff --git a/src/algorithm/REDUCE_SUM-Hip.cpp b/src/algorithm/REDUCE_SUM-Hip.cpp
@@ -266,7 +266,8 @@ void REDUCE_SUM::runHipVariantRAJANewReduce(VariantID vid)
       RAJA::forall<exec_policy>( res,
         RAJA::RangeSegment(ibegin, iend),
         RAJA::expt::Reduce<RAJA::operators::plus>(&tsum),
-        [=] __device__ (Index_type i, Real_type& sum) {
+        [=] __device__ (Index_type i,
+          RAJA::expt::ValOp<Real_type, RAJA::operators::plus>& sum) {
           REDUCE_SUM_BODY;
         }
       );

diff --git a/src/algorithm/REDUCE_SUM-OMP.cpp b/src/algorithm/REDUCE_SUM-OMP.cpp
@@ -76,14 +76,16 @@ void REDUCE_SUM::runOpenMPVariant(VariantID vid, size_t tune_idx)
 
     case RAJA_OpenMP : {
 
+      RAJA::resources::Host res;
+
       if (tune_idx == 0) {
 
         startTimer();
         for (RepIndex_type irep = 0; irep < run_reps; ++irep) {
 
           RAJA::ReduceSum<RAJA::omp_reduce, Real_type> sum(m_sum_init);
 
-          RAJA::forall<RAJA::omp_parallel_for_exec>(
+          RAJA::forall<RAJA::omp_parallel_for_exec>(res,
             RAJA::RangeSegment(ibegin, iend),
             [=](Index_type i) {
               REDUCE_SUM_BODY;
@@ -101,10 +103,11 @@ void REDUCE_SUM::runOpenMPVariant(VariantID vid, size_t tune_idx)
 
           Real_type tsum = m_sum_init;
 
-          RAJA::forall<RAJA::omp_parallel_for_exec>(
+          RAJA::forall<RAJA::omp_parallel_for_exec>(res,
             RAJA::RangeSegment(ibegin, iend),
             RAJA::expt::Reduce<RAJA::operators::plus>(&tsum),
-            [=] (Index_type i, Real_type& sum) {
+            [=] (Index_type i,
+              RAJA::expt::ValOp<Real_type, RAJA::operators::plus>& sum) {
               REDUCE_SUM_BODY;
             }
           );

diff --git a/src/algorithm/REDUCE_SUM-OMPTarget.cpp b/src/algorithm/REDUCE_SUM-OMPTarget.cpp
@@ -64,7 +64,8 @@ void REDUCE_SUM::runOpenMPTargetVariant(VariantID vid, size_t RAJAPERF_UNUSED_AR
       RAJA::forall<RAJA::omp_target_parallel_for_exec<threads_per_team>>(
         RAJA::RangeSegment(ibegin, iend),
         RAJA::expt::Reduce<RAJA::operators::plus>(&tsum),
-        [=] (Index_type i, Real_type& sum) {
+        [=] (Index_type i,
+          RAJA::expt::ValOp<Real_type, RAJA::operators::plus>& sum) {
           REDUCE_SUM_BODY;
         }
       );

diff --git a/src/algorithm/REDUCE_SUM-Seq.cpp b/src/algorithm/REDUCE_SUM-Seq.cpp
@@ -76,15 +76,17 @@ void REDUCE_SUM::runSeqVariant(VariantID vid, size_t tune_idx)
 
     case RAJA_Seq : {
 
+      RAJA::resources::Host res;
+
       if (tune_idx == 0) {
 
         startTimer();
         for (RepIndex_type irep = 0; irep < run_reps; ++irep) {
 
           RAJA::ReduceSum<RAJA::seq_reduce, Real_type> sum(m_sum_init);
 
-          RAJA::forall<RAJA::seq_exec>( RAJA::RangeSegment(ibegin, iend),
-            [=](Index_type i) {
+          RAJA::forall<RAJA::seq_exec>(res,
+            RAJA::RangeSegment(ibegin, iend), [=](Index_type i) {
               REDUCE_SUM_BODY;
           });
 
@@ -100,9 +102,11 @@ void REDUCE_SUM::runSeqVariant(VariantID vid, size_t tune_idx)
 
           Real_type tsum = m_sum_init;
 
-          RAJA::forall<RAJA::seq_exec>( RAJA::RangeSegment(ibegin, iend),
+          RAJA::forall<RAJA::seq_exec>(res,
+            RAJA::RangeSegment(ibegin, iend),
             RAJA::expt::Reduce<RAJA::operators::plus>(&tsum),
-            [=] (Index_type i, Real_type& sum) {
+            [=] (Index_type i,
+              RAJA::expt::ValOp<Real_type, RAJA::operators::plus>& sum) {
               REDUCE_SUM_BODY;
             }
           );

diff --git a/src/algorithm/REDUCE_SUM-Sycl.cpp b/src/algorithm/REDUCE_SUM-Sycl.cpp
@@ -76,11 +76,13 @@ void REDUCE_SUM::runSyclVariantImpl(VariantID vid)
     for (RepIndex_type irep = 0; irep < run_reps; ++irep) {
 
        Real_type tsum = m_sum_init;
+
        RAJA::forall< RAJA::sycl_exec<work_group_size, true /*async*/> >( 
          res,
          RAJA::RangeSegment(ibegin, iend), 
          RAJA::expt::Reduce<RAJA::operators::plus>(&tsum),
-         [=]  (Index_type i, Real_type& sum) {
+         [=]  (Index_type i,
+           RAJA::expt::ValOp<Real_type, RAJA::operators::plus>& sum) {
            REDUCE_SUM_BODY;
          }
        );

diff --git a/src/basic/PI_REDUCE-Cuda.cpp b/src/basic/PI_REDUCE-Cuda.cpp
@@ -168,7 +168,8 @@ void PI_REDUCE::runCudaVariantRAJANewReduce(VariantID vid)
       RAJA::forall< exec_policy >( res,
         RAJA::RangeSegment(ibegin, iend),
         RAJA::expt::Reduce<RAJA::operators::plus>(&tpi),
-        [=] __device__ (Index_type i, Real_type& pi) {
+        [=] __device__ (Index_type i,
+          RAJA::expt::ValOp<Real_type, RAJA::operators::plus>& pi) {
           PI_REDUCE_BODY;
         }
       );

diff --git a/src/basic/PI_REDUCE-Hip.cpp b/src/basic/PI_REDUCE-Hip.cpp
@@ -168,7 +168,8 @@ void PI_REDUCE::runHipVariantRAJANewReduce(VariantID vid)
         res,
         RAJA::RangeSegment(ibegin, iend),
         RAJA::expt::Reduce<RAJA::operators::plus>(&tpi),
-        [=] __device__ (Index_type i, Real_type& pi) {
+        [=] __device__ (Index_type i,
+          RAJA::expt::ValOp<Real_type, RAJA::operators::plus>& pi) {
           PI_REDUCE_BODY;
         }
       );

diff --git a/src/basic/PI_REDUCE-OMP.cpp b/src/basic/PI_REDUCE-OMP.cpp
@@ -77,16 +77,17 @@ void PI_REDUCE::runOpenMPVariant(VariantID vid, size_t tune_idx)
 
     case RAJA_OpenMP : {
 
+      RAJA::resources::Host res;
+
       if (tune_idx == 0) {
 
         startTimer();
         for (RepIndex_type irep = 0; irep < run_reps; ++irep) {
 
           RAJA::ReduceSum<RAJA::omp_reduce, Real_type> pi(m_pi_init);
 
-          RAJA::forall<RAJA::omp_parallel_for_exec>(
-            RAJA::RangeSegment(ibegin, iend),
-            [=](Index_type i) {
+          RAJA::forall<RAJA::omp_parallel_for_exec>(res,
+            RAJA::RangeSegment(ibegin, iend), [=](Index_type i) {
               PI_REDUCE_BODY;
           });
 
@@ -102,10 +103,11 @@ void PI_REDUCE::runOpenMPVariant(VariantID vid, size_t tune_idx)
 
           Real_type tpi = m_pi_init;
 
-          RAJA::forall<RAJA::omp_parallel_for_exec>(
+          RAJA::forall<RAJA::omp_parallel_for_exec>(res,
             RAJA::RangeSegment(ibegin, iend),
             RAJA::expt::Reduce<RAJA::operators::plus>(&tpi),
-            [=] (Index_type i, Real_type& pi) {
+            [=] (Index_type i,
+              RAJA::expt::ValOp<Real_type, RAJA::operators::plus>& pi) {
               PI_REDUCE_BODY;
             }
           );

diff --git a/src/basic/PI_REDUCE-OMPTarget.cpp b/src/basic/PI_REDUCE-OMPTarget.cpp
@@ -64,7 +64,8 @@ void PI_REDUCE::runOpenMPTargetVariant(VariantID vid, size_t RAJAPERF_UNUSED_ARG
       RAJA::forall<RAJA::omp_target_parallel_for_exec<threads_per_team>>(
         RAJA::RangeSegment(ibegin, iend),
         RAJA::expt::Reduce<RAJA::operators::plus>(&tpi),
-        [=] (Index_type i, Real_type& pi) {
+        [=] (Index_type i,
+          RAJA::expt::ValOp<Real_type, RAJA::operators::plus>& pi) {
           PI_REDUCE_BODY;
         }
       );