Updates camp submodule

.github/workflows/build.yml

@@ -54,6 +54,7 @@ jobs:
 ## Config and build action
       - uses: threeal/cmake-action@v1.3.0
         with:
+          cmake-version: '3.29.x'  # Back off version of cmake with known issue
           build-dir: build
           options:
             ENABLE_WARNINGS_AS_ERRORS=Off

.uberenv_config.json

@@ -6,6 +6,6 @@
 "spack_url": "https://github.com/spack/spack.git",
 "spack_branch": "v1.1.1",
 "spack_configs_path": "scripts/radiuss-spack-configs",
-"spack_packages_path": "scripts/radiuss-spack-configs/spack_repo/llnl_radiuss/packages",
+"spack_packages_path": "spack_repo/raja/packages",
 "spack_setup_clingo": false
 }

examples/memoryManager.hpp

@@ -77,8 +77,8 @@ void deallocate(T *&ptr)
 #elif defined(RAJA_ENABLE_HIP)
     CAMP_HIP_API_INVOKE_AND_CHECK(hipMalloc, (void **)&ptr, sizeof(T) * size);
 #elif defined(RAJA_ENABLE_SYCL)
-      auto qu = sycl_res->get<camp::resources::Sycl>().get_queue();
-      ptr = ::sycl::malloc_device<T>(size, *qu);
+      auto& qu = sycl_res->get<camp::resources::Sycl>().get_queue();
+      ptr = ::sycl::malloc_device<T>(size, qu);
 #endif
     return ptr;
   }

examples/resource-forall.cpp

@@ -271,7 +271,7 @@ using EXEC_POLICY = RAJA::sycl_exec<GPU_BLOCK_SIZE>;
   // _raja_res_k2_end
 
   // _raja_res_wait_start
-  res_gpu2.wait_for(&e);
+  res_gpu2.wait_for(e);
   // _raja_res_wait_end
 
   // _raja_res_k3_start

examples/resource-kernel.cpp

@@ -56,7 +56,7 @@ int main(int RAJA_UNUSED_ARG(argc), char **RAJA_UNUSED_ARG(argv[]))
         }
       );
 
-      def_cuda_res.wait_for(&e);
+      def_cuda_res.wait_for(e);
     }
   );
 

examples/resource-launch.cpp

@@ -57,7 +57,7 @@ int main(int RAJA_UNUSED_ARG(argc), char **RAJA_UNUSED_ARG(argv[]))
 
       });
 
-      def_cuda_res.wait_for(&e);
+      def_cuda_res.wait_for(e);
     }
   );
 

exercises/memoryManager.hpp

@@ -77,8 +77,8 @@ void deallocate(T *&ptr)
 #elif defined(RAJA_ENABLE_HIP)
     CAMP_HIP_API_INVOKE_AND_CHECK(hipMalloc, (void **)&ptr, sizeof(T) * size);
 #elif defined(RAJA_ENABLE_SYCL)
-      auto qu = sycl_res->get<camp::resources::Sycl>().get_queue();
-      ptr = ::sycl::malloc_device<T>(size, *qu);
+      auto& qu = sycl_res->get<camp::resources::Sycl>().get_queue();
+      ptr = ::sycl::malloc_device<T>(size, qu);
 #endif
     return ptr;
   }

include/RAJA/pattern/concepts.hpp

@@ -110,7 +110,12 @@ RAJAMakeExecPolWithIterMappingConcept(DirectBasePolicy,
 
 namespace type_traits
 {
-DefineTypeTraitFromConcept(is_execution_policy, concepts::ExecutionPolicy);
+template<typename T>
+struct is_execution_policy : std::bool_constant<concepts::ExecutionPolicy<T>>
+{};
+
+template<typename T>
+inline constexpr bool is_execution_policy_v = is_execution_policy<T>::value;
 
 }  // namespace type_traits
 

include/RAJA/pattern/kernel/For.hpp

@@ -110,7 +110,7 @@ struct StatementExecutor<
     auto len    = segment_length<ArgumentId>(data);
     using len_t = decltype(len);
 
-    auto r = data.res;
+    auto&& r = data.get_resource();
 
     forall_impl(r, ExecPolicy {}, TypedRangeSegment<len_t>(0, len), for_wrapper,
                 RAJA::expt::get_empty_forall_param_pack());

include/RAJA/pattern/kernel/internal/LoopData.hpp

@@ -129,7 +129,7 @@ struct LoopData
       typename RAJA::expt::detail::ParamToArgHelper<ParamTuple>::type;
   ParamTuple param_tuple;
 
-  Resource res;
+  Resource* res;
 
   // Lambdas that were passed into the kernel
   using BodiesTuple = camp::tuple<Bodies...>;
@@ -143,11 +143,11 @@ struct LoopData
 
   RAJA_INLINE RAJA_HOST_DEVICE constexpr LoopData(SegmentTuple const& s,
                                                   ParamTuple const& p,
-                                                  Resource r,
+                                                  Resource& r,
                                                   Bodies const&... b)
       : segment_tuple(s),
         param_tuple(p),
-        res(r),
+        res(&r),
         bodies(b...)
   {}
 
@@ -175,7 +175,7 @@ struct LoopData
     return camp::get<ParamId::param_idx>(param_tuple);
   }
 
-  RAJA_HOST_DEVICE RAJA_INLINE Resource get_resource() { return res; }
+  RAJA_INLINE Resource& get_resource() { return *res; }
 };
 
 template<camp::idx_t ArgumentId, typename Data>

include/RAJA/policy/sycl/MemUtils_SYCL.hpp

@@ -66,6 +66,10 @@ extern std::unordered_map<::sycl::queue, bool> g_queue_info_map;
 
 }  // namespace detail
 
+//! get queue for current launch
+RAJA_INLINE
+::sycl::queue& currentResourceQueue() { return detail::tl_status.qu; }
+
 //! Allocator for pinned memory for use in basic_mempool
 struct PinnedAllocator
 {
@@ -74,17 +78,17 @@ struct PinnedAllocator
   void* malloc(size_t nbytes)
   {
     void* ptr;
-    ::sycl::queue* q = ::camp::resources::Sycl::get_default().get_queue();
-    ptr              = ::sycl::malloc_host(nbytes, *q);
+    ::sycl::queue& q = ::camp::resources::Sycl::get_default().get_queue();
+    ptr              = ::sycl::malloc_host(nbytes, q);
     return ptr;
   }
 
   // returns true on success
   // Will throw if ptr is not in q's context
   bool free(void* ptr)
   {
-    ::sycl::queue* q = ::camp::resources::Sycl::get_default().get_queue();
-    ::sycl::free(ptr, *q);
+    ::sycl::queue& q = ::camp::resources::Sycl::get_default().get_queue();
+    ::sycl::free(ptr, q);
     return true;
   }
 };
@@ -97,17 +101,17 @@ struct DeviceAllocator
   void* malloc(size_t nbytes)
   {
     void* ptr;
-    ::sycl::queue* q = ::camp::resources::Sycl::get_default().get_queue();
-    ptr              = ::sycl::malloc_device(nbytes, *q);
+    ::sycl::queue& q = ::camp::resources::Sycl::get_default().get_queue();
+    ptr              = ::sycl::malloc_device(nbytes, q);
     return ptr;
   }
 
   // returns true on success
   // Will throw if ptr is not in q's context
   bool free(void* ptr)
   {
-    ::sycl::queue* q = ::camp::resources::Sycl::get_default().get_queue();
-    ::sycl::free(ptr, *q);
+    ::sycl::queue& q = ::camp::resources::Sycl::get_default().get_queue();
+    ::sycl::free(ptr, q);
     return true;
   }
 };
@@ -121,18 +125,18 @@ struct DeviceZeroedAllocator
   void* malloc(size_t nbytes)
   {
     void* ptr;
-    ::sycl::queue* q = ::camp::resources::Sycl::get_default().get_queue();
-    ptr              = ::sycl::malloc_device(nbytes, *q);
-    q->memset(ptr, 0, nbytes);
+    ::sycl::queue& q = ::camp::resources::Sycl::get_default().get_queue();
+    ptr              = ::sycl::malloc_device(nbytes, q);
+    q.memset(ptr, 0, nbytes);
     return ptr;
   }
 
   // Returns true on success
   // Will throw if ptr is not in q's context
   bool free(void* ptr)
   {
-    ::sycl::queue* q = ::camp::resources::Sycl::get_default().get_queue();
-    ::sycl::free(ptr, *q);
+    ::sycl::queue& q = ::camp::resources::Sycl::get_default().get_queue();
+    ::sycl::free(ptr, q);
     return true;
   }
 };

include/RAJA/policy/sycl/forall.hpp

@@ -130,7 +130,7 @@ RAJA_INLINE resources::EventProxy<resources::Sycl> forall_impl(
   sycl_dim_t blockSize {BlockSize};
   sycl_dim_t gridSize = impl::getGridDim(static_cast<size_t>(len), BlockSize);
 
-  ::sycl::queue* q  = sycl_res.get_queue();
+  ::sycl::queue& q  = sycl_res.get_queue();
   LOOP_BODY* lbody  = nullptr;
   Iterator* d_begin = nullptr;
 
@@ -145,11 +145,11 @@ RAJA_INLINE resources::EventProxy<resources::Sycl> forall_impl(
     // Kernel body is nontrivially copyable, create space on device and copy to
     // Workaround until "is_device_copyable" is supported
     //
-    lbody = (LOOP_BODY*)::sycl::malloc_device(sizeof(LoopBody), *q);
-    q->memcpy(lbody, &loop_body, sizeof(LOOP_BODY)).wait();
+    lbody = (LOOP_BODY*)::sycl::malloc_device(sizeof(LoopBody), q);
+    q.memcpy(lbody, &loop_body, sizeof(LOOP_BODY)).wait();
 
-    d_begin = (Iterator*)::sycl::malloc_device(sizeof(Iterator), *q);
-    q->memcpy(d_begin, &begin, sizeof(Iterator)).wait();
+    d_begin = (Iterator*)::sycl::malloc_device(sizeof(Iterator), q);
+    q.memcpy(d_begin, &begin, sizeof(Iterator)).wait();
   }
 
   // Both the parallel_for call, combinations, and resolution are all
@@ -161,11 +161,11 @@ RAJA_INLINE resources::EventProxy<resources::Sycl> forall_impl(
       return x;
     };
 
-    ForallParam* res = ::sycl::malloc_shared<ForallParam>(1, *q);
+    ForallParam* res = ::sycl::malloc_shared<ForallParam>(1, q);
     RAJA::expt::ParamMultiplexer::parampack_init(pol, *res);
     auto reduction = ::sycl::reduction(res, f_params, combiner);
 
-    q->submit([&](::sycl::handler& h) {
+    q.submit([&](::sycl::handler& h) {
       h.parallel_for(::sycl::range<1>(len), reduction,
                      [=](::sycl::item<1> it, auto& red) {
                        ForallParam fp;
@@ -186,15 +186,15 @@ RAJA_INLINE resources::EventProxy<resources::Sycl> forall_impl(
                      });
     });
 
-    q->wait();
+    q.wait();
     RAJA::expt::ParamMultiplexer::parampack_combine(pol, f_params, *res);
-    ::sycl::free(res, *q);
+    ::sycl::free(res, q);
     RAJA::expt::ParamMultiplexer::parampack_resolve(pol, f_params);
   }
   // Note: separate branches
   else
   {
-    q->submit([&](::sycl::handler& h) {
+    q.submit([&](::sycl::handler& h) {
       h.parallel_for(::sycl::nd_range<1> {gridSize, blockSize},
                      [=](::sycl::nd_item<1> it) {
                        IndexType ii = it.get_global_id(0);
@@ -214,16 +214,16 @@ RAJA_INLINE resources::EventProxy<resources::Sycl> forall_impl(
 
     if (!Async)
     {
-      q->wait();
+      q.wait();
     }
   }
 
 
   // If we had to allocate device memory, free it
   if constexpr (!is_lbody_trivially_copyable)
   {
-    ::sycl::free(lbody, *q);
-    ::sycl::free(d_begin, *q);
+    ::sycl::free(lbody, q);
+    ::sycl::free(d_begin, q);
   }
 
 
@@ -267,8 +267,8 @@ RAJA_INLINE resources::EventProxy<resources::Sycl> forall_impl(
 
   if (!Async)
   {
-    ::sycl::queue* q = r.get_queue();
-    q->wait();
+    ::sycl::queue& q = r.get_queue();
+    q.wait();
   }
 
   return resources::EventProxy<resources::Sycl>(r);

include/RAJA/policy/sycl/kernel/SyclKernel.hpp

@@ -213,7 +213,7 @@ struct StatementExecutor<
                                       LaunchConfig, stmt_list_t, data_t, Types>;
 
     camp::resources::Sycl res = data.get_resource();
-    ::sycl::queue* q          = res.get_queue();
+    ::sycl::queue& q          = res.get_queue();
     ;
 
     //
@@ -226,7 +226,7 @@ struct StatementExecutor<
     //
     // Launch the kernels
     //
-    launch_t::launch(std::move(data), launch_dims, shmem, q);
+    launch_t::launch(std::move(data), launch_dims, shmem, &q);
   }
 };
 

include/RAJA/policy/sycl/launch.hpp

@@ -54,7 +54,7 @@ struct LaunchExecute<RAJA::sycl_launch_t<async, 0>>
     EXEC_POL pol {};
 
     /*Get the queue from concrete resource */
-    ::sycl::queue* q = res.get<camp::resources::Sycl>().get_queue();
+    ::sycl::queue& q = res.get<camp::resources::Sycl>().get_queue();
 
     if constexpr (!is_parampack_empty)
     {
@@ -94,8 +94,8 @@ struct LaunchExecute<RAJA::sycl_launch_t<async, 0>>
     //
     if constexpr (!is_lbody_trivially_copyable)
     {
-      lbody = (LOOP_BODY*)::sycl::malloc_device(sizeof(LOOP_BODY), *q);
-      q->memcpy(lbody, &loop_body, sizeof(LOOP_BODY)).wait();
+      lbody = (LOOP_BODY*)::sycl::malloc_device(sizeof(LOOP_BODY), q);
+      q.memcpy(lbody, &loop_body, sizeof(LOOP_BODY)).wait();
     }
     // Both the parallel_for call, combinations, and resolution are all
     // unique to the parameter case, so we make a constexpr branch here
@@ -106,11 +106,11 @@ struct LaunchExecute<RAJA::sycl_launch_t<async, 0>>
         return x;
       };
 
-      ReduceParams* res = ::sycl::malloc_shared<ReduceParams>(1, *q);
+      ReduceParams* res = ::sycl::malloc_shared<ReduceParams>(1, q);
       RAJA::expt::ParamMultiplexer::parampack_init(pol, *res);
       auto reduction = ::sycl::reduction(res, launch_reducers, combiner);
 
-      q->submit([&](::sycl::handler& h) {
+      q.submit([&](::sycl::handler& h) {
          auto s_vec =
              ::sycl::local_accessor<char, 1>(launch_params.shared_mem_size, h);
 
@@ -141,13 +141,13 @@ struct LaunchExecute<RAJA::sycl_launch_t<async, 0>>
 
       RAJA::expt::ParamMultiplexer::parampack_combine(pol, launch_reducers,
                                                       *res);
-      ::sycl::free(res, *q);
-      ::sycl::free(lbody, *q);
+      ::sycl::free(res, q);
+      ::sycl::free(lbody, q);
       RAJA::expt::ParamMultiplexer::parampack_resolve(pol, launch_reducers);
     }
     else
     {
-      q->submit([&](::sycl::handler& h) {
+      q.submit([&](::sycl::handler& h) {
         auto s_vec =
             ::sycl::local_accessor<char, 1>(launch_params.shared_mem_size, h);
 
@@ -173,7 +173,7 @@ struct LaunchExecute<RAJA::sycl_launch_t<async, 0>>
 
       if (!async)
       {
-        q->wait();
+        q.wait();
       }
     }