From b895bdfdb1027653eae46de2b4c5f779dfa4332f Mon Sep 17 00:00:00 2001 From: Brandon Echols Date: Mon, 13 Jul 2026 14:09:59 -0700 Subject: [PATCH 01/15] Updates camp submodule --- tpl/camp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tpl/camp b/tpl/camp index e75ab64c02..824a6a3ba4 160000 --- a/tpl/camp +++ b/tpl/camp @@ -1 +1 @@ -Subproject commit e75ab64c029aa27c80593715cb2a3ccad7453c8c +Subproject commit 824a6a3ba48a233791332c398e7f024b2191aa27 From 9320d51f526d7e39f6b6a2a28bbeb00f870770a0 Mon Sep 17 00:00:00 2001 From: Brandon Echols Date: Mon, 13 Jul 2026 14:12:00 -0700 Subject: [PATCH 02/15] Converts wait_for event parameters to references --- examples/resource-forall.cpp | 2 +- examples/resource-kernel.cpp | 2 +- examples/resource-launch.cpp | 2 +- test/unit/resource/tests/test-resource-Depends.hpp | 2 +- test/unit/resource/tests/test-resource-JoinAsyncSemantics.hpp | 4 ++-- test/unit/resource/tests/test-resource-MultiStream.hpp | 4 ++-- 6 files changed, 8 insertions(+), 8 deletions(-) diff --git a/examples/resource-forall.cpp b/examples/resource-forall.cpp index 26a88c6e8a..885e1705c6 100644 --- a/examples/resource-forall.cpp +++ b/examples/resource-forall.cpp @@ -271,7 +271,7 @@ using EXEC_POLICY = RAJA::sycl_exec; // _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 diff --git a/examples/resource-kernel.cpp b/examples/resource-kernel.cpp index e32d66f5a7..f5b9314aad 100644 --- a/examples/resource-kernel.cpp +++ b/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); } ); diff --git a/examples/resource-launch.cpp b/examples/resource-launch.cpp index 840f2a38a6..3ad636daf1 100644 --- a/examples/resource-launch.cpp +++ b/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); } ); diff --git a/test/unit/resource/tests/test-resource-Depends.hpp b/test/unit/resource/tests/test-resource-Depends.hpp index 14bdc0ace3..c18f044435 100644 --- a/test/unit/resource/tests/test-resource-Depends.hpp +++ b/test/unit/resource/tests/test-resource-Depends.hpp @@ -39,7 +39,7 @@ void ResourceDependsTestImpl() } ); - dev1.wait_for(&e); + dev1.wait_for(e); forall(dev1, RangeSegment(0,ARRAY_SIZE), [=] RAJA_HOST_DEVICE (int i) { diff --git a/test/unit/resource/tests/test-resource-JoinAsyncSemantics.hpp b/test/unit/resource/tests/test-resource-JoinAsyncSemantics.hpp index 79129db55b..b30595f82c 100644 --- a/test/unit/resource/tests/test-resource-JoinAsyncSemantics.hpp +++ b/test/unit/resource/tests/test-resource-JoinAsyncSemantics.hpp @@ -34,7 +34,7 @@ void ResourceJoinAsyncSemanticsTestImpl() dev2.memcpy(d_array, h_array, sizeof(int) * ARRAY_SIZE); auto e1 = dev2.get_event_erased(); - dev1.wait_for(&e1); + dev1.wait_for(e1); RAJA::resources::Event e2 = forall(dev1, RangeSegment(0,ARRAY_SIZE), [=] RAJA_HOST_DEVICE (int i) { @@ -42,7 +42,7 @@ void ResourceJoinAsyncSemanticsTestImpl() } ); - dev2.wait_for(&e2); + dev2.wait_for(e2); dev2.memcpy(h_array, d_array, sizeof(int) * ARRAY_SIZE); diff --git a/test/unit/resource/tests/test-resource-MultiStream.hpp b/test/unit/resource/tests/test-resource-MultiStream.hpp index fcc87e3b40..9247663dd9 100644 --- a/test/unit/resource/tests/test-resource-MultiStream.hpp +++ b/test/unit/resource/tests/test-resource-MultiStream.hpp @@ -47,8 +47,8 @@ void ResourceMultiStreamTestImpl() } }); - dev1.wait_for(&e2); - dev1.wait_for(&e3); + dev1.wait_for(e2); + dev1.wait_for(e3); dev1.memcpy(h_array, d_array, sizeof(int) * ARRAY_SIZE); From 27c74378d80f7720b661dfc621d32274507591b1 Mon Sep 17 00:00:00 2001 From: Brandon Echols Date: Mon, 13 Jul 2026 14:22:47 -0700 Subject: [PATCH 03/15] Fixes to SYCL build --- include/RAJA/policy/sycl/MemUtils_SYCL.hpp | 24 ++++++------ include/RAJA/policy/sycl/forall.hpp | 30 +++++++-------- .../RAJA/policy/sycl/kernel/SyclKernel.hpp | 4 +- include/RAJA/policy/sycl/launch.hpp | 18 ++++----- include/RAJA/policy/sycl/reduce.hpp | 38 +++++++------------ 5 files changed, 51 insertions(+), 63 deletions(-) diff --git a/include/RAJA/policy/sycl/MemUtils_SYCL.hpp b/include/RAJA/policy/sycl/MemUtils_SYCL.hpp index 13c941779f..ea5ac056c7 100644 --- a/include/RAJA/policy/sycl/MemUtils_SYCL.hpp +++ b/include/RAJA/policy/sycl/MemUtils_SYCL.hpp @@ -74,8 +74,8 @@ 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; } @@ -83,8 +83,8 @@ struct PinnedAllocator // 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,8 +97,8 @@ 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; } @@ -106,8 +106,8 @@ struct DeviceAllocator // 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,8 +121,8 @@ 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); + ::sycl::queue& q = ::camp::resources::Sycl::get_default().get_queue(); + ptr = ::sycl::malloc_device(nbytes, q); q->memset(ptr, 0, nbytes); return ptr; } @@ -131,8 +131,8 @@ struct DeviceZeroedAllocator // 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; } }; diff --git a/include/RAJA/policy/sycl/forall.hpp b/include/RAJA/policy/sycl/forall.hpp index 7d886c79c1..b6c4106dff 100644 --- a/include/RAJA/policy/sycl/forall.hpp +++ b/include/RAJA/policy/sycl/forall.hpp @@ -130,7 +130,7 @@ RAJA_INLINE resources::EventProxy forall_impl( sycl_dim_t blockSize {BlockSize}; sycl_dim_t gridSize = impl::getGridDim(static_cast(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 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 forall_impl( return x; }; - ForallParam* res = ::sycl::malloc_shared(1, *q); + ForallParam* res = ::sycl::malloc_shared(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 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,7 +214,7 @@ RAJA_INLINE resources::EventProxy forall_impl( if (!Async) { - q->wait(); + q.wait(); } } @@ -222,8 +222,8 @@ RAJA_INLINE resources::EventProxy forall_impl( // 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 forall_impl( if (!Async) { - ::sycl::queue* q = r.get_queue(); - q->wait(); + ::sycl::queue& q = r.get_queue(); + q.wait(); } return resources::EventProxy(r); diff --git a/include/RAJA/policy/sycl/kernel/SyclKernel.hpp b/include/RAJA/policy/sycl/kernel/SyclKernel.hpp index c31aca7ee8..72192963a9 100644 --- a/include/RAJA/policy/sycl/kernel/SyclKernel.hpp +++ b/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); } }; diff --git a/include/RAJA/policy/sycl/launch.hpp b/include/RAJA/policy/sycl/launch.hpp index 02b7a53c84..0f210e44a1 100644 --- a/include/RAJA/policy/sycl/launch.hpp +++ b/include/RAJA/policy/sycl/launch.hpp @@ -54,7 +54,7 @@ struct LaunchExecute> EXEC_POL pol {}; /*Get the queue from concrete resource */ - ::sycl::queue* q = res.get().get_queue(); + ::sycl::queue& q = res.get().get_queue(); if constexpr (!is_parampack_empty) { @@ -94,8 +94,8 @@ struct LaunchExecute> // 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> return x; }; - ReduceParams* res = ::sycl::malloc_shared(1, *q); + ReduceParams* res = ::sycl::malloc_shared(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(launch_params.shared_mem_size, h); @@ -141,13 +141,13 @@ struct LaunchExecute> 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(launch_params.shared_mem_size, h); @@ -173,7 +173,7 @@ struct LaunchExecute> if (!async) { - q->wait(); + q.wait(); } } diff --git a/include/RAJA/policy/sycl/reduce.hpp b/include/RAJA/policy/sycl/reduce.hpp index a3cd200528..3fe4a8f0ce 100644 --- a/include/RAJA/policy/sycl/reduce.hpp +++ b/include/RAJA/policy/sycl/reduce.hpp @@ -120,13 +120,13 @@ struct Reduce_Data Reduce_Data(T initValue, T identityValue, Offload_Info& info) : value(initValue) { - ::sycl::queue* q = ::camp::resources::Sycl::get_default().get_queue(); + ::sycl::queue& q = ::camp::resources::Sycl::get_default().get_queue(); device = reinterpret_cast( - ::sycl::malloc_device(sycl::MaxNumTeams * sizeof(T), *(q))); + ::sycl::malloc_device(sycl::MaxNumTeams * sizeof(T), q)); host = reinterpret_cast( - ::sycl::malloc_host(sycl::MaxNumTeams * sizeof(T), *(q))); + ::sycl::malloc_host(sycl::MaxNumTeams * sizeof(T), q)); if (!host) { @@ -150,18 +150,12 @@ struct Reduce_Data //! transfers from the host to the device -- exit() is called upon failure RAJA_INLINE void hostToDevice(Offload_Info& RAJA_UNUSED_ARG(info)) { - ::sycl::queue* q = ::camp::resources::Sycl::get_default().get_queue(); - - if (!q) - { - camp::resources::Resource res = camp::resources::Sycl(); - q = res.get().get_queue(); - } + ::sycl::queue& q = ::camp::resources::Sycl::get_default().get_queue(); // precondition: host and device are valid pointers auto e = - q->memcpy(reinterpret_cast(device), - reinterpret_cast(host), sycl::MaxNumTeams * sizeof(T)); + q.memcpy(reinterpret_cast(device), + reinterpret_cast(host), sycl::MaxNumTeams * sizeof(T)); e.wait(); } @@ -169,18 +163,12 @@ struct Reduce_Data //! transfers from the device to the host -- exit() is called upon failure RAJA_INLINE void deviceToHost(Offload_Info& RAJA_UNUSED_ARG(info)) { - ::sycl::queue* q = ::camp::resources::Sycl::get_default().get_queue(); - - if (!q) - { - camp::resources::Resource res = camp::resources::Sycl(); - q = res.get().get_queue(); - } + ::sycl::queue& q = ::camp::resources::Sycl::get_default().get_queue(); // precondition: host and device are valid pointers - auto e = q->memcpy(reinterpret_cast(host), - reinterpret_cast(device), - sycl::MaxNumTeams * sizeof(T)); + auto e = q.memcpy(reinterpret_cast(host), + reinterpret_cast(device), + sycl::MaxNumTeams * sizeof(T)); e.wait(); } @@ -188,16 +176,16 @@ struct Reduce_Data //! frees all data from the offload information passed RAJA_INLINE void cleanup(Offload_Info& RAJA_UNUSED_ARG(info)) { - ::sycl::queue* q = ::camp::resources::Sycl::get_default().get_queue(); + ::sycl::queue& q = ::camp::resources::Sycl::get_default().get_queue(); if (device) { - ::sycl::free(reinterpret_cast(device), *q); + ::sycl::free(reinterpret_cast(device), q); device = nullptr; } if (host) { - ::sycl::free(reinterpret_cast(host), *q); + ::sycl::free(reinterpret_cast(host), q); // delete[] host; host = nullptr; } From afc62bcf0e64b31375b8f11c482beda3a5db2088 Mon Sep 17 00:00:00 2001 From: Brandon Echols Date: Mon, 13 Jul 2026 14:23:46 -0700 Subject: [PATCH 04/15] Formats code --- include/RAJA/policy/sycl/reduce.hpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/include/RAJA/policy/sycl/reduce.hpp b/include/RAJA/policy/sycl/reduce.hpp index 3fe4a8f0ce..0ed51d222d 100644 --- a/include/RAJA/policy/sycl/reduce.hpp +++ b/include/RAJA/policy/sycl/reduce.hpp @@ -154,8 +154,8 @@ struct Reduce_Data // precondition: host and device are valid pointers auto e = - q.memcpy(reinterpret_cast(device), - reinterpret_cast(host), sycl::MaxNumTeams * sizeof(T)); + q.memcpy(reinterpret_cast(device), reinterpret_cast(host), + sycl::MaxNumTeams * sizeof(T)); e.wait(); } @@ -166,9 +166,9 @@ struct Reduce_Data ::sycl::queue& q = ::camp::resources::Sycl::get_default().get_queue(); // precondition: host and device are valid pointers - auto e = q.memcpy(reinterpret_cast(host), - reinterpret_cast(device), - sycl::MaxNumTeams * sizeof(T)); + auto e = + q.memcpy(reinterpret_cast(host), reinterpret_cast(device), + sycl::MaxNumTeams * sizeof(T)); e.wait(); } From 81b19265a0ad198c71966eb7c909a8f6ded28205 Mon Sep 17 00:00:00 2001 From: Brandon Echols Date: Mon, 13 Jul 2026 14:35:44 -0700 Subject: [PATCH 05/15] Adds copy operator --- include/RAJA/policy/sycl/reduce.hpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/include/RAJA/policy/sycl/reduce.hpp b/include/RAJA/policy/sycl/reduce.hpp index 0ed51d222d..26528af129 100644 --- a/include/RAJA/policy/sycl/reduce.hpp +++ b/include/RAJA/policy/sycl/reduce.hpp @@ -147,6 +147,9 @@ struct Reduce_Data //! default copy constructor for POD Reduce_Data(const Reduce_Data&) = default; + //! default copy operator for POD + Reduce_Data& operator=(const Reduce_Data&) = default; + //! transfers from the host to the device -- exit() is called upon failure RAJA_INLINE void hostToDevice(Offload_Info& RAJA_UNUSED_ARG(info)) { From 9e592bfa394db523f9895f8b575ef6cf93709b50 Mon Sep 17 00:00:00 2001 From: Brandon Echols Date: Mon, 13 Jul 2026 15:26:37 -0700 Subject: [PATCH 06/15] Fixes accessor --- include/RAJA/policy/sycl/MemUtils_SYCL.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/RAJA/policy/sycl/MemUtils_SYCL.hpp b/include/RAJA/policy/sycl/MemUtils_SYCL.hpp index ea5ac056c7..fc100449dd 100644 --- a/include/RAJA/policy/sycl/MemUtils_SYCL.hpp +++ b/include/RAJA/policy/sycl/MemUtils_SYCL.hpp @@ -123,7 +123,7 @@ struct DeviceZeroedAllocator void* ptr; ::sycl::queue& q = ::camp::resources::Sycl::get_default().get_queue(); ptr = ::sycl::malloc_device(nbytes, q); - q->memset(ptr, 0, nbytes); + q.memset(ptr, 0, nbytes); return ptr; } From d5271388f234a05fdcf776b2e6999a23572730e6 Mon Sep 17 00:00:00 2001 From: Brandon Echols Date: Mon, 13 Jul 2026 15:29:17 -0700 Subject: [PATCH 07/15] Updates concept due to camp update --- include/RAJA/pattern/concepts.hpp | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/include/RAJA/pattern/concepts.hpp b/include/RAJA/pattern/concepts.hpp index 2514eb3478..47c19f1fb0 100644 --- a/include/RAJA/pattern/concepts.hpp +++ b/include/RAJA/pattern/concepts.hpp @@ -110,7 +110,12 @@ RAJAMakeExecPolWithIterMappingConcept(DirectBasePolicy, namespace type_traits { -DefineTypeTraitFromConcept(is_execution_policy, concepts::ExecutionPolicy); +template +struct is_execution_policy : std::bool_constant> +{}; + +template +inline constexpr bool is_execution_policy_v = is_execution_policy::value; } // namespace type_traits From 8102e6fd315bc93b991e672488601583d6ac6667 Mon Sep 17 00:00:00 2001 From: Brandon Echols Date: Mon, 13 Jul 2026 17:12:41 -0700 Subject: [PATCH 08/15] Moves to use references instead of pointers due to camp update --- examples/memoryManager.hpp | 4 ++-- exercises/memoryManager.hpp | 4 ++-- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/examples/memoryManager.hpp b/examples/memoryManager.hpp index 7afdbeafc0..aeea7c26ed 100644 --- a/examples/memoryManager.hpp +++ b/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().get_queue(); - ptr = ::sycl::malloc_device(size, *qu); + auto& qu = sycl_res->get().get_queue(); + ptr = ::sycl::malloc_device(size, qu); #endif return ptr; } diff --git a/exercises/memoryManager.hpp b/exercises/memoryManager.hpp index 7afdbeafc0..aeea7c26ed 100644 --- a/exercises/memoryManager.hpp +++ b/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().get_queue(); - ptr = ::sycl::malloc_device(size, *qu); + auto& qu = sycl_res->get().get_queue(); + ptr = ::sycl::malloc_device(size, qu); #endif return ptr; } From 7c2b4a59805f2a585fcc990dab2ef067eb17fd02 Mon Sep 17 00:00:00 2001 From: Jason Burmark Date: Tue, 14 Jul 2026 10:09:33 -0700 Subject: [PATCH 09/15] Resource by pointer in kernel LoopData This is fine as the LoopData object is created on the stack in a function where the resource is an argument. Note that the LoopData object is used on the device in async kernels but that is also fine as the resource may not be used on the device. --- include/RAJA/pattern/kernel/For.hpp | 2 +- include/RAJA/pattern/kernel/internal/LoopData.hpp | 8 ++++---- 2 files changed, 5 insertions(+), 5 deletions(-) diff --git a/include/RAJA/pattern/kernel/For.hpp b/include/RAJA/pattern/kernel/For.hpp index 572da59fef..ebfa680fa0 100644 --- a/include/RAJA/pattern/kernel/For.hpp +++ b/include/RAJA/pattern/kernel/For.hpp @@ -110,7 +110,7 @@ struct StatementExecutor< auto len = segment_length(data); using len_t = decltype(len); - auto r = data.res; + auto&& r = data.get_resource(); forall_impl(r, ExecPolicy {}, TypedRangeSegment(0, len), for_wrapper, RAJA::expt::get_empty_forall_param_pack()); diff --git a/include/RAJA/pattern/kernel/internal/LoopData.hpp b/include/RAJA/pattern/kernel/internal/LoopData.hpp index ad9aea0cfc..07e2c0a312 100644 --- a/include/RAJA/pattern/kernel/internal/LoopData.hpp +++ b/include/RAJA/pattern/kernel/internal/LoopData.hpp @@ -129,7 +129,7 @@ struct LoopData typename RAJA::expt::detail::ParamToArgHelper::type; ParamTuple param_tuple; - Resource res; + Resource* res; // Lambdas that were passed into the kernel using BodiesTuple = camp::tuple; @@ -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(param_tuple); } - RAJA_HOST_DEVICE RAJA_INLINE Resource get_resource() { return res; } + RAJA_INLINE Resource& get_resource() { return *res; } }; template From 216c39542b619e590e7ee409c2023f40e9b8dc34 Mon Sep 17 00:00:00 2001 From: Jason Burmark Date: Tue, 14 Jul 2026 10:10:04 -0700 Subject: [PATCH 10/15] Update clang format version in amd/cray scripts --- scripts/lc-builds/toss4_amdclang.sh | 2 +- scripts/lc-builds/toss4_amdclang_asan.sh | 2 +- scripts/lc-builds/toss4_amdclang_proteus.sh | 2 +- scripts/lc-builds/toss4_cce_hip.sh | 2 +- scripts/lc-builds/toss4_hipcc.sh | 2 +- 5 files changed, 5 insertions(+), 5 deletions(-) diff --git a/scripts/lc-builds/toss4_amdclang.sh b/scripts/lc-builds/toss4_amdclang.sh index f61a8fe307..4196fb88a2 100755 --- a/scripts/lc-builds/toss4_amdclang.sh +++ b/scripts/lc-builds/toss4_amdclang.sh @@ -80,7 +80,7 @@ cmake \ -DHIP_ROOT_DIR="/opt/rocm-${COMP_VER}/hip" \ -DHIP_PATH=/opt/rocm-${COMP_VER}/llvm/bin \ -DENABLE_CLANGFORMAT=On \ - -DCLANGFORMAT_EXECUTABLE=/opt/rocm-5.2.3/llvm/bin/clang-format \ + -DCLANGFORMAT_EXECUTABLE=/opt/rocm-6.4.3/llvm/bin/clang-format \ -DCMAKE_C_COMPILER=/opt/rocm-${COMP_VER}/llvm/bin/amdclang \ -DCMAKE_CXX_COMPILER=/opt/rocm-${COMP_VER}/llvm/bin/amdclang++ \ -DCMAKE_HIP_ARCHITECTURES="${COMP_ARCH}" \ diff --git a/scripts/lc-builds/toss4_amdclang_asan.sh b/scripts/lc-builds/toss4_amdclang_asan.sh index 527b60c0ca..059747ddd9 100755 --- a/scripts/lc-builds/toss4_amdclang_asan.sh +++ b/scripts/lc-builds/toss4_amdclang_asan.sh @@ -95,7 +95,7 @@ cmake \ -DCMAKE_C_FLAGS="-fsanitize=address -shared-libsan" \ -DCMAKE_CXX_FLAGS="-fsanitize=address -shared-libsan" \ -DENABLE_CLANGFORMAT=On \ - -DCLANGFORMAT_EXECUTABLE=/opt/rocm-5.2.3/llvm/bin/clang-format \ + -DCLANGFORMAT_EXECUTABLE=/opt/rocm-6.4.3/llvm/bin/clang-format \ -DBLT_CXX_STD=c++20 \ -C "../host-configs/lc-builds/toss4/${HOSTCONFIG}.cmake" \ -DENABLE_HIP=ON \ diff --git a/scripts/lc-builds/toss4_amdclang_proteus.sh b/scripts/lc-builds/toss4_amdclang_proteus.sh index c7ec284a4a..913dbfb9d9 100755 --- a/scripts/lc-builds/toss4_amdclang_proteus.sh +++ b/scripts/lc-builds/toss4_amdclang_proteus.sh @@ -80,7 +80,7 @@ cmake \ -DHIP_ROOT_DIR="/opt/rocm-${COMP_VER}/hip" \ -DHIP_PATH=/opt/rocm-${COMP_VER}/llvm/bin \ -DENABLE_CLANGFORMAT=On \ - -DCLANGFORMAT_EXECUTABLE=/opt/rocm-5.2.3/llvm/bin/clang-format \ + -DCLANGFORMAT_EXECUTABLE=/opt/rocm-6.4.3/llvm/bin/clang-format \ -DCMAKE_C_COMPILER=/opt/rocm-${COMP_VER}/llvm/bin/amdclang \ -DCMAKE_CXX_COMPILER=/opt/rocm-${COMP_VER}/llvm/bin/amdclang++ \ -DCMAKE_HIP_ARCHITECTURES="${COMP_ARCH}" \ diff --git a/scripts/lc-builds/toss4_cce_hip.sh b/scripts/lc-builds/toss4_cce_hip.sh index 3863608770..27e00e8fba 100755 --- a/scripts/lc-builds/toss4_cce_hip.sh +++ b/scripts/lc-builds/toss4_cce_hip.sh @@ -70,7 +70,7 @@ cmake \ -DAMDGPU_TARGETS=${HIP_ARCH} \ -DBLT_CXX_STD=c++20 \ -DENABLE_CLANGFORMAT=On \ - -DCLANGFORMAT_EXECUTABLE=/opt/rocm-5.2.3/llvm/bin/clang-format \ + -DCLANGFORMAT_EXECUTABLE=/opt/rocm-6.4.3/llvm/bin/clang-format \ -C "../host-configs/lc-builds/toss4/${HOSTCONFIG}.cmake" \ -DENABLE_HIP=ON \ -DENABLE_OPENMP=ON \ diff --git a/scripts/lc-builds/toss4_hipcc.sh b/scripts/lc-builds/toss4_hipcc.sh index 22592f7a8f..c81dec7c7d 100755 --- a/scripts/lc-builds/toss4_hipcc.sh +++ b/scripts/lc-builds/toss4_hipcc.sh @@ -87,7 +87,7 @@ cmake \ -DAMDGPU_TARGETS="${COMP_ARCH}" \ -DBLT_CXX_STD=c++20 \ -DENABLE_CLANGFORMAT=On \ - -DCLANGFORMAT_EXECUTABLE=/opt/rocm-5.2.3/llvm/bin/clang-format \ + -DCLANGFORMAT_EXECUTABLE=/opt/rocm-6.4.3/llvm/bin/clang-format \ -C "../host-configs/lc-builds/toss4/${HOSTCONFIG}.cmake" \ -DENABLE_HIP=ON \ -DENABLE_OPENMP=ON \ From fd993292b75d50ae66f799c88bae49f2d8807326 Mon Sep 17 00:00:00 2001 From: Brandon Echols Date: Wed, 15 Jul 2026 12:54:56 -0700 Subject: [PATCH 11/15] Adds temporary package files for raja CI --- .uberenv_config.json | 2 +- .../raja/packages/camp/camp-rocm6.patch | 15 + .../camp/libstdc++-13-missing-header.patch | 24 + spack_repo/raja/packages/camp/package.py | 162 ++++ spack_repo/raja/packages/raja/package.py | 754 ++++++++++++++++++ ...tile-iterator-comparison-fix-2024.02.patch | 20 + spack_repo/raja/repo.yaml | 3 + 7 files changed, 979 insertions(+), 1 deletion(-) create mode 100644 spack_repo/raja/packages/camp/camp-rocm6.patch create mode 100644 spack_repo/raja/packages/camp/libstdc++-13-missing-header.patch create mode 100644 spack_repo/raja/packages/camp/package.py create mode 100644 spack_repo/raja/packages/raja/package.py create mode 100644 spack_repo/raja/packages/raja/tile-iterator-comparison-fix-2024.02.patch create mode 100644 spack_repo/raja/repo.yaml diff --git a/.uberenv_config.json b/.uberenv_config.json index 452aa08baf..e5904922f1 100644 --- a/.uberenv_config.json +++ b/.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 } diff --git a/spack_repo/raja/packages/camp/camp-rocm6.patch b/spack_repo/raja/packages/camp/camp-rocm6.patch new file mode 100644 index 0000000000..dbb3fa0220 --- /dev/null +++ b/spack_repo/raja/packages/camp/camp-rocm6.patch @@ -0,0 +1,15 @@ +diff -ruN spack-src/include/camp/resource/hip.hpp spack-src-patched/include/camp/resource/hip.hpp +--- spack-src/include/camp/resource/hip.hpp 2021-08-20 23:38:39.000000000 +0000 ++++ spack-src-patched/include/camp/resource/hip.hpp 2024-06-06 21:34:25.174477941 +0000 +@@ -111,7 +111,11 @@ + hipPointerAttribute_t a; + hipError_t status = hipPointerGetAttributes(&a, p); + if (status == hipSuccess) { ++#if (HIP_VERSION_MAJOR >= 6) ++ switch (a.type) { ++#else + switch (a.memoryType) { ++#endif + case hipMemoryTypeHost: + return MemoryAccess::Pinned; + case hipMemoryTypeDevice: diff --git a/spack_repo/raja/packages/camp/libstdc++-13-missing-header.patch b/spack_repo/raja/packages/camp/libstdc++-13-missing-header.patch new file mode 100644 index 0000000000..d3c642abe6 --- /dev/null +++ b/spack_repo/raja/packages/camp/libstdc++-13-missing-header.patch @@ -0,0 +1,24 @@ +diff --git a/include/camp/resource.hpp b/include/camp/resource.hpp +index 58df206..763aa20 100644 +--- a/include/camp/resource.hpp ++++ b/include/camp/resource.hpp +@@ -12,6 +12,7 @@ http://github.com/llnl/camp + #define __CAMP_RESOURCE_HPP + + #include ++#include + #include + #include + #include +diff --git a/include/camp/resource/host.hpp b/include/camp/resource/host.hpp +index c765175..4428c57 100644 +--- a/include/camp/resource/host.hpp ++++ b/include/camp/resource/host.hpp +@@ -13,6 +13,7 @@ http://github.com/llnl/camp + + #include "camp/resource/event.hpp" + #include "camp/resource/platform.hpp" ++#include + + namespace camp + { diff --git a/spack_repo/raja/packages/camp/package.py b/spack_repo/raja/packages/camp/package.py new file mode 100644 index 0000000000..e1ec7f9783 --- /dev/null +++ b/spack_repo/raja/packages/camp/package.py @@ -0,0 +1,162 @@ +# Copyright Spack Project Developers. See COPYRIGHT file for details. +# +# SPDX-License-Identifier: (Apache-2.0 OR MIT) + +from spack_repo.builtin.build_systems.cached_cmake import cmake_cache_string +from spack_repo.builtin.build_systems.cmake import CMakePackage +from spack_repo.builtin.build_systems.cuda import CudaPackage +from spack_repo.builtin.build_systems.rocm import ROCmPackage + +from spack.package import * + + +class Camp(CMakePackage, CudaPackage, ROCmPackage): + """ + Compiler agnostic metaprogramming library providing concepts, + type operations and tuples for C++ and cuda + """ + + homepage = "https://github.com/LLNL/camp" + git = "https://github.com/LLNL/camp.git" + url = "https://github.com/LLNL/camp/archive/v0.1.0.tar.gz" + + maintainers("adrienbernede", "kab163", "trws") + + license("BSD-3-Clause") + + version("main", branch="main", submodules=False) + version( + "2026.07.0", + tag="v2026.07.0", + commit="824a6a3ba48a233791332c398e7f024b2191aa27", + submodules=False, + ) + version( + "2025.12.0", + tag="v2025.12.0", + commit="a8caefa9f4c811b1a114b4ed2c9b681d40f12325", + submodules=False, + ) + version( + "2025.09.2", + tag="v2025.09.2", + commit="4070ce93a802849d61037310a87c50cc24c9e498", + submodules=False, + ) + version( + "2025.09.0", + tag="v2025.09.0", + commit="b642f29b9d0eee9113bea2791958c29243063e5c", + submodules=False, + ) + version( + "2025.03.0", + tag="v2025.03.0", + commit="ee0a3069a7ae72da8bcea63c06260fad34901d43", + submodules=False, + ) + version( + "2024.07.0", + tag="v2024.07.0", + commit="0f07de4240c42e0b38a8d872a20440cb4b33d9f5", + submodules=False, + ) + version( + "2024.02.1", + tag="v2024.02.1", + commit="79c320fa09db987923b56884afdc9f82f4b70fc4", + submodules=False, + ) + version( + "2024.02.0", + tag="v2024.02.0", + commit="03c80a6c6ab4f97e76a52639563daec71435a277", + submodules=False, + ) + version( + "2023.06.0", + tag="v2023.06.0", + commit="ac34c25b722a06b138bc045d38bfa5e8fa3ec9c5", + submodules=False, + ) + version("2022.10.1", sha256="2d12f1a46f5a6d01880fc075cfbd332e2cf296816a7c1aa12d4ee5644d386f02") + version("2022.10.0", sha256="3561c3ef00bbcb61fe3183c53d49b110e54910f47e7fc689ad9ccce57e55d6b8") + version("2022.03.2", sha256="bc4aaeacfe8f2912e28f7a36fc731ab9e481bee15f2c6daf0cb208eed3f201eb") + version("2022.03.0", sha256="e9090d5ee191ea3a8e36b47a8fe78f3ac95d51804f1d986d931e85b8f8dad721") + version("0.3.0", sha256="129431a049ca5825443038ad5a37a86ba6d09b2618d5fe65d35f83136575afdb") + version("0.2.3", sha256="58a0f3bd5eadb588d7dc83f3d050aff8c8db639fc89e8d6553f9ce34fc2421a7") + version("0.2.2", sha256="194d38b57e50e3494482a7f94940b27f37a2bee8291f2574d64db342b981d819") + version("0.1.0", sha256="fd4f0f2a60b82a12a1d9f943f8893dc6fe770db493f8fae5ef6f7d0c439bebcc") + + # TODO: figure out gtest dependency and then set this default True. + variant("tests", default=False, description="Build tests") + variant("openmp", default=False, description="Build with OpenMP support") + variant("omptarget", default=False, description="Build with OpenMP Target support") + variant("sycl", default=False, description="Build with Sycl support") + + depends_on("c", type="build") + depends_on("cxx", type="build") + + with when("+cuda"): + depends_on("cub", when="^cuda@:10") + + depends_on("blt", type="build") + depends_on("blt@0.7.1:", type="build", when="@2025.09.0:") + depends_on("blt@0.7.0:", type="build", when="@2025.03.0:") + depends_on("blt@0.6.2:", type="build", when="@2024.02.1:") + depends_on("blt@0.6.1", type="build", when="@2024.02.0") + depends_on("blt@0.5.0:0.5.3", type="build", when="@2022.03.0:2023.06.0") + + patch("libstdc++-13-missing-header.patch", when="@:2022.10") + + patch("camp-rocm6.patch", when="@0.2.3 +rocm ^hip@6:") + + conflicts("^blt@:0.3.6", when="+rocm") + + conflicts("+omptarget +rocm") + conflicts("+sycl +omptarget") + conflicts("+sycl +rocm") + conflicts( + "+sycl", + when="@:2024.02.99", + msg="Support for SYCL was introduced in RAJA after 2024.02 release, " + "please use a newer release.", + ) + + def cmake_args(self): + spec = self.spec + + options = [] + + options.append("-DBLT_SOURCE_DIR={0}".format(spec["blt"].prefix)) + + options.append(self.define_from_variant("ENABLE_CUDA", "cuda")) + if spec.satisfies("+cuda"): + options.append("-DCUDA_TOOLKIT_ROOT_DIR={0}".format(spec["cuda"].prefix)) + + if not spec.satisfies("cuda_arch=none"): + cuda_arch = spec.variants["cuda_arch"].value + options.append("-DCMAKE_CUDA_ARCHITECTURES={0}".format(cuda_arch[0])) + options.append("-DCUDA_ARCH=sm_{0}".format(cuda_arch[0])) + flag = "-arch sm_{0}".format(cuda_arch[0]) + options.append("-DCMAKE_CUDA_FLAGS:STRING={0}".format(flag)) + + options.append(self.define_from_variant("ENABLE_HIP", "rocm")) + if spec.satisfies("+rocm"): + rocm_root = spec["llvm-amdgpu"].prefix + options.append(self.define("ROCM_PATH", rocm_root)) + + archs = ";".join(self.spec.variants["amdgpu_target"].value) + options.append("-DCMAKE_HIP_ARCHITECTURES={0}".format(archs)) + options.append("-DGPU_TARGETS={0}".format(archs)) + options.append("-DAMDGPU_TARGETS={0}".format(archs)) + + if spec.satisfies("+omptarget"): + options.append(cmake_cache_string("RAJA_DATA_ALIGN", 64)) + + options.append(self.define_from_variant("ENABLE_TESTS", "tests")) + options.append(self.define_from_variant("ENABLE_OPENMP", "openmp")) + options.append(self.define_from_variant("CAMP_ENABLE_TARGET_OPENMP", "omptarget")) + options.append(self.define_from_variant("ENABLE_SYCL", "sycl")) + + return options diff --git a/spack_repo/raja/packages/raja/package.py b/spack_repo/raja/packages/raja/package.py new file mode 100644 index 0000000000..3300f1e694 --- /dev/null +++ b/spack_repo/raja/packages/raja/package.py @@ -0,0 +1,754 @@ +# Copyright Spack Project Developers. See COPYRIGHT file for details. +# +# SPDX-License-Identifier: (Apache-2.0 OR MIT) + +import os +import re +import shutil +import socket +from textwrap import dedent + +from spack_repo.builtin.build_systems.cached_cmake import ( + CachedCMakePackage, + cmake_cache_option, + cmake_cache_path, + cmake_cache_string, +) +from spack_repo.builtin.build_systems.cuda import CudaPackage +from spack_repo.builtin.build_systems.rocm import ROCmPackage +from spack_repo.builtin.packages.blt.package import llnl_link_helpers + +from spack.package import * + + +# Starting with 2022.03.0, the only submodule we want to fetch is tpl/desul +# since there is no package for it. Other RAJA submodules are defined as +# dependencies. +def submodules(package): + submodules = [] + submodules.append("tpl/desul") + return submodules + + +class Raja(CachedCMakePackage, CudaPackage, ROCmPackage): + """RAJA Parallel Framework.""" + + homepage = "https://github.com/LLNL/RAJA" + git = "https://github.com/LLNL/RAJA.git" + tags = ["radiuss", "e4s"] + + maintainers("adrienbernede", "davidbeckingsale", "kab163") + + license("BSD-3-Clause") + + version("develop", branch="develop", submodules=submodules) + version("main", branch="main", submodules=submodules) + version( + "2025.12.2", + tag="v2025.12.2", + commit="eca7c5015a5cf8bf7cc8ad1829fd36d3276ab274", + submodules=submodules, + ) + version( + "2025.12.1", + tag="v2025.12.1", + commit="3b8b59a1e9be2e1066c0d77372b3bf5956e6d6e2", + submodules=submodules, + ) + version( + "2025.12.0", + tag="v2025.12.0", + commit="e827035c630e71a9358e2f21c2f3cf6fd5fb6605", + submodules=submodules, + ) + version( + "2025.09.1", + tag="v2025.09.1", + commit="1e0756eda3c344da362e483afb9100ebd8137a2c", + submodules=submodules, + ) + version( + "2025.09.0", + tag="v2025.09.0", + commit="ca756788dbdd43fec2a3840389126ae94a905d5f", + submodules=submodules, + ) + version( + "2025.03.2", + tag="v2025.03.2", + commit="6e36a94380adbe88fed11a3213fc08461428ece0", + submodules=submodules, + ) + version( + "2025.03.1", + tag="v2025.03.1", + commit="ffa7b92377705aff855b4bf602e197ae4f8e8cc3", + submodules=submodules, + ) + version( + "2025.03.0", + tag="v2025.03.0", + commit="1d70abf171474d331f1409908bdf1b1c3fe19222", + submodules=submodules, + ) + version( + "2024.07.0", + tag="v2024.07.0", + commit="4d7fcba55ebc7cb972b7cc9f6778b48e43792ea1", + submodules=submodules, + ) + version( + "2024.02.2", + tag="v2024.02.2", + commit="593f756b14ac57ded33ee61d8d2292d4beb840e6", + submodules=submodules, + ) + version( + "2024.02.1", + tag="v2024.02.1", + commit="3ada0950b0774ec907d30a9eceaf6af7478b833b", + submodules=submodules, + ) + version( + "2024.02.0", + tag="v2024.02.0", + commit="82d1b926ada0fbb15a4a6e0adadc30c715cfda7b", + submodules=submodules, + ) + version( + "2023.06.1", + tag="v2023.06.1", + commit="9b5f61edf3aa1e6fdbc9a4b30828c81504639963", + submodules=submodules, + ) + version( + "2023.06.0", + tag="v2023.06.0", + commit="e330b2560747d5417cd7bd265fab3fb91d32ecbd", + submodules=submodules, + ) + version( + "2022.10.5", + tag="v2022.10.5", + commit="3774f51339459bbbdb77055aa23f82919b6335b6", + submodules=submodules, + ) + version( + "2022.10.4", + tag="v2022.10.4", + commit="c2a6b1740759ae3ae7c85b35e20dbffbe235355d", + submodules=submodules, + ) + version( + "2022.03.0", + tag="v2022.03.0", + commit="4351fe6a50bd579511a625b017c9e054885e7fd2", + submodules=submodules, + ) + version( + "0.14.0", tag="v0.14.0", commit="357933a42842dd91de5c1034204d937fce0a2a44", submodules=True + ) + version( + "0.13.0", tag="v0.13.0", commit="3047fa720132d19ee143b1fcdacaa72971f5988c", submodules=True + ) + version( + "0.12.1", tag="v0.12.1", commit="9cb6370bb2868e35ebba23cdce927f5f7f9da530", submodules=True + ) + version( + "0.12.0", tag="v0.12.0", commit="32d92e38da41cc8d4db25ec79b9884a73a0cb3a1", submodules=True + ) + version( + "0.11.0", tag="v0.11.0", commit="0502b9b69c4cb60aa0afbdf699b555c76cb18f22", submodules=True + ) + version( + "0.10.1", tag="v0.10.1", commit="be91e040130678b1350dbda56cc352433db758bd", submodules=True + ) + version( + "0.10.0", tag="v0.10.0", commit="53cb89cf788d28bc4ed2b4e6f75483fdd26024aa", submodules=True + ) + version( + "0.9.0", tag="v0.9.0", commit="df7ca1fa892b6ac4147c614d2d739d5022f63fc7", submodules=True + ) + version( + "0.8.0", tag="v0.8.0", commit="8d19a8c2cbac611de6f92ad8852b9f3454b27e63", submodules=True + ) + version( + "0.7.0", tag="v0.7.0", commit="caa33b371b586dfae3d8569caee91c5eddfd7b31", submodules=True + ) + version( + "0.6.0", tag="v0.6.0", commit="cc7a97e8b4e52c3de820c9dfacd358822a147871", submodules=True + ) + version( + "0.5.3", tag="v0.5.3", commit="1ca35c0ed2a43a3fa9c6cd70c5d25f16d88ecd8c", submodules=True + ) + version( + "0.5.2", tag="v0.5.2", commit="4d5c3d5d7f311838855f7010810610349e729f64", submodules=True + ) + version( + "0.5.1", tag="v0.5.1", commit="bf340abe5199d7e051520913c9a7a5de336b5820", submodules=True + ) + version( + "0.5.0", tag="v0.5.0", commit="9b539d84fdad049f65caeba836f41031f5baf4cc", submodules=True + ) + version( + "0.4.1", tag="v0.4.1", commit="3618cfe95d6a442fa50fbe7bfbcf654cf9f800b9", submodules=True + ) + version( + "0.4.0", tag="v0.4.0", commit="31b2a48192542c2da426885baa5af0ed57606b78", submodules=True + ) + + # export targets when building pre-2.4.0 release with BLT 0.4.0+ + patch( + "https://github.com/LLNL/RAJA/commit/eca1124ee4af380d6613adc6012c307d1fd4176b.patch?full_index=1", + sha256="12bb78c00b6683ad3e7fd4e3f87f9776bae074b722431b79696bc862816735ef", + when="@:0.13.0 ^blt@0.4:", + ) + + # Backward compatibility is stopped from ROCm 6.0 + # Future relase will have the change from PR https://github.com/LLNL/RAJA/pull/1568 + patch( + "https://github.com/LLNL/RAJA/commit/406eb8dee05a41eb32c421c375688a4863b60642.patch?full_index=1", + sha256="d9ce5ef038555cbccb330a9016b7be77e56ae0660583cba955dab9d0297a4b07", + when="^hip@6.0", + ) + + # Fix compilation issue reported by Intel from their new compiler version + patch( + "https://github.com/LLNL/RAJA/commit/3e831e034bd92daacf49f40b66459aefd6ea3972.patch?full_index=1", + sha256="c0548fc5220f24082fb2592d5b4e8b7c8c783b87906d5f0950d53953d25161f6", + when="@2024.02.1:2024.02.99 %oneapi@2025:", + ) + + patch("tile-iterator-comparison-fix-2024.02.patch", when="@2024.02.0:2024.02.2") + + variant("openmp", default=False, description="Build OpenMP backend") + variant("shared", default=False, description="Build shared libs") + variant("desul", default=False, description="Build desul atomics backend") + variant("vectorization", default=True, description="Build SIMD/SIMT intrinsics support") + variant( + "omptask", default=False, description="Build OpenMP task variants of internal algorithms" + ) + variant("omptarget", default=False, description="Build OpenMP on target device support") + variant("sycl", default=False, description="Build sycl backend") + variant("gpu-profiling", default=False, description="Enable GPU profiling") + + variant("plugins", default=False, description="Enable runtime plugins") + variant("caliper", default=False, description="Enable caliper support") + variant("examples", default=True, description="Build examples.") + variant("exercises", default=True, description="Build exercises.") + # TODO: figure out gtest dependency and then set this default True + # and remove the +tests conflict below. + variant("tests", default=False, description="Build tests") + + # we don't use variants to express the failing test, we only add a variant to + # define whether we want to run all the tests (including those known to fail) + # or only the passing ones. + variant( + "run-all-tests", + default=False, + description="Run all the tests, including those known to fail.", + ) + + variant( + "lowopttest", + default=False, + description="For developers, lowers optimization level to pass tests with some compilers", + ) + + variant( + "cxxstd", + default="20", + values=("11", "14", "17", "20"), + description="C++ standard to build with", + ) + conflicts("cxxstd=11", when="@0.14.0:") + conflicts("cxxstd=14", when="@2025.09.0:") + conflicts("cxxstd=17", when="@2026.03.0:") + conflicts("+sycl cxxstd=14", when="@2024.07.0:") + + depends_on("cxx", type="build") + depends_on("c", type="build") + + depends_on("blt", type="build") + # TODO(smith84): Edit the following line after the June 2026 RAJA suite release + depends_on("blt@0.7.2:", type="build", when="@develop") + depends_on("blt@0.7.1:", type="build", when="@2025.09.0:") + depends_on("blt@0.7.0:", type="build", when="@2025.03.0:") + depends_on("blt@0.6.2:", type="build", when="@2024.02.1:") + depends_on("blt@0.6.1", type="build", when="@2024.02.0") + depends_on("blt@0.5.3", type="build", when="@2023.06.0:2023.06.1") + depends_on("blt@0.5.2:0.5.3", type="build", when="@2022.10.5") + depends_on("blt@0.5.0:0.5.3", type="build", when="@0.14.1:2022.10.4") + depends_on("blt@0.4.1", type="build", when="@0.14.0") + depends_on("blt@0.4.0:0.4.1", type="build", when="@0.13.0") + depends_on("blt@0.3.6:0.4.1", type="build", when="@:0.12.0") + conflicts("^blt@:0.3.6", when="+rocm") + conflicts("^blt@:0.7.1", when="+cuda ^cuda@13:", msg="CUDA 13+ requires BLT 0.7.2 or newer") + + depends_on("camp") + depends_on("camp+openmp", when="+openmp") + depends_on("camp+omptarget", when="+omptarget") + depends_on("camp+sycl", when="+sycl") + # TODO(johnbowen42): Remove the following line after the June 2026 RAJA suite release + depends_on("camp@2026.07", when="@develop") + depends_on("camp@2025.12", when="@2025.12.0:2025.12.2") + depends_on("camp@2025.09", when="@2025.09") + depends_on("camp@2025.03", when="@2025.03") + depends_on("camp@2024.07", when="@2024.07") + depends_on("camp@2024.02.1", when="@2024.02.1") + depends_on("camp@2024.02.0", when="@2024.02.0") + depends_on("camp@2023.06.0", when="@2023.06.0:2023.06.1") + depends_on("camp@2022.10.1:2023.06.0", when="@2022.10.3:2022.10.5") + depends_on("camp@2022.10.0:2023.06.0", when="@2022.10.0:2022.10.2") + depends_on("camp@2022.03.2", when="@2022.03.0:2022.03.1") + depends_on("camp@0.2.2:0.2.3", when="@0.14.0") + depends_on("camp@0.1.0", when="@0.10.0:0.13.0") + + depends_on("cmake@3.24:", when="@2025.09.0:", type="build") + depends_on("cmake@3.23:", when="@2024.07.0:2025.03.2", type="build") + depends_on("cmake@3.23:", when="@2022.10.0:2024.02.2+rocm", type="build") + depends_on("cmake@3.20:", when="@2022.10.0:2024.02.2", type="build") + depends_on("cmake@3.20:", when="@:2022.03+rocm", type="build") + depends_on("cmake@3.14:", when="@:2022.03", type="build") + + depends_on("llvm-openmp", when="+openmp %apple-clang") + + depends_on("caliper", when="+caliper") + + depends_on("rocprim", when="+rocm") + with when("+rocm @0.12.0:"): + depends_on("camp+rocm") + for arch in ROCmPackage.amdgpu_targets: + depends_on( + "camp+rocm amdgpu_target={0}".format(arch), when="amdgpu_target={0}".format(arch) + ) + conflicts("+openmp", when="@:2022.03") + + with when("+cuda @0.12.0:"): + depends_on("camp+cuda") + for sm_ in CudaPackage.cuda_arch_values: + depends_on("camp +cuda cuda_arch={0}".format(sm_), when="cuda_arch={0}".format(sm_)) + + conflicts("+gpu-profiling", when="~cuda~rocm", msg="GPU profiling requires CUDA or ROCm") + conflicts("+gpu-profiling +cuda", when="@:2022.02.99") + conflicts("+gpu-profiling +rocm", when="@:2022.02.99") + + conflicts("+omptarget +rocm") + conflicts("+sycl +omptarget") + conflicts("+sycl +rocm") + conflicts( + "+sycl", + when="@:2024.02.99", + msg="Support for SYCL was introduced in RAJA after 2024.02 release, " + "please use a newer release.", + ) + + depends_on("cuda@12:", when="+cuda") + conflicts( + "^cuda@13:", + when="@:2025.12.2 +cuda", + msg="RAJA versions up to and including 2025.12.2 do not support CUDA 13+", + ) + + def _get_sys_type(self, spec): + sys_type = spec.architecture + if "SYS_TYPE" in env: + sys_type = env["SYS_TYPE"] + return sys_type + + @property + def libs(self): + shared = "+shared" in self.spec + return find_libraries("libRAJA", root=self.prefix, shared=shared, recursive=True) + + @property + def cache_name(self): + hostname = socket.gethostname() + if "SYS_TYPE" in env: + hostname = hostname.rstrip("1234567890") + return "{0}-{1}-{2}@{3}-{4}.cmake".format( + hostname, + self._get_sys_type(self.spec), + self.spec.compiler.name, + self.spec.compiler.version, + self.spec.dag_hash(8), + ) + + def initconfig_compiler_entries(self): + spec = self.spec + compiler = self.compiler + # Default entries are already defined in CachedCMakePackage, inherit them: + entries = super().initconfig_compiler_entries() + + if spec.satisfies("+rocm ^blt@:0.6"): + entries.insert(0, cmake_cache_path("CMAKE_CXX_COMPILER", spec["hip"].hipcc)) + + llnl_link_helpers(entries, spec, compiler) + + return entries + + def initconfig_hardware_entries(self): + spec = self.spec + entries = super().initconfig_hardware_entries() + + entries.append("#------------------{0}".format("-" * 30)) + entries.append("# Package custom hardware settings") + entries.append("#------------------{0}\n".format("-" * 30)) + + entries.append(cmake_cache_option("ENABLE_OPENMP", spec.satisfies("+openmp"))) + entries.append(cmake_cache_option("ENABLE_CUDA", spec.satisfies("+cuda"))) + + if spec.satisfies("+cuda"): + # CUDA configuration from cuda_for_radiuss_projects + cuda_flags = [] + if not spec.satisfies("cuda_arch=none"): + cuda_archs = ";".join(spec.variants["cuda_arch"].value) + entries.append(cmake_cache_string("CMAKE_CUDA_ARCHITECTURES", cuda_archs)) + + # gcc-toolchain support + gcc_toolchain_regex = re.compile(".*gcc-toolchain.*") + using_toolchain = list( + filter(gcc_toolchain_regex.match, spec.compiler_flags["cxxflags"]) + ) + if using_toolchain: + cuda_flags.append("-Xcompiler {}".format(using_toolchain[0])) + + if cuda_flags: + entries.append(cmake_cache_string("CMAKE_CUDA_FLAGS", " ".join(cuda_flags))) + + if spec.satisfies("+rocm"): + entries.append(cmake_cache_option("ENABLE_HIP", True)) + + # HIP configuration from hip_for_radiuss_projects + rocm_root = spec["llvm-amdgpu"].prefix + gcc_toolchain_regex = re.compile(".*gcc-toolchain.*") + using_toolchain = list( + filter(gcc_toolchain_regex.match, spec.compiler_flags["cxxflags"]) + ) + hip_link_flags = "" + + if using_toolchain: + gcc_prefix = using_toolchain[0] + entries.append( + cmake_cache_string("HIP_CLANG_FLAGS", "--gcc-toolchain={0}".format(gcc_prefix)) + ) + entries.append( + cmake_cache_string( + "CMAKE_EXE_LINKER_FLAGS", + hip_link_flags + " -Wl,-rpath={0}/lib64".format(gcc_prefix), + ) + ) + else: + entries.append( + cmake_cache_string( + "CMAKE_EXE_LINKER_FLAGS", "-Wl,-rpath={0}/llvm/lib/".format(rocm_root) + ) + ) + + hipcc_flags = [] + if self.spec.satisfies("^rocprim@7.0"): + hipcc_flags.append("-std=c++17") + if self.spec.satisfies("@2025.09.0:"): + hipcc_flags.append("-std=c++17") + elif self.spec.satisfies("@0.14.0:2025.09.0"): + hipcc_flags.append("-std=c++14") + entries.append(cmake_cache_string("HIP_HIPCC_FLAGS", " ".join(hipcc_flags))) + else: + entries.append(cmake_cache_option("ENABLE_HIP", False)) + + return entries + + @property + def cxx_std(self): + return self.spec.variants.get("cxxstd").value + + def initconfig_package_entries(self): + spec = self.spec + entries = [] + + option_prefix = "RAJA_" if spec.satisfies("@0.14.0:") else "" + + # TPL locations + entries.append("#------------------{0}".format("-" * 60)) + entries.append("# TPLs") + entries.append("#------------------{0}\n".format("-" * 60)) + + entries.append(cmake_cache_path("BLT_SOURCE_DIR", spec["blt"].prefix)) + if "camp" in self.spec: + entries.append(cmake_cache_path("camp_DIR", spec["camp"].prefix)) + + # Build options + entries.append("#------------------{0}".format("-" * 60)) + entries.append("# Build Options") + entries.append("#------------------{0}\n".format("-" * 60)) + + entries.append(cmake_cache_string("CMAKE_BUILD_TYPE", spec.variants["build_type"].value)) + entries.append(cmake_cache_option("BUILD_SHARED_LIBS", spec.satisfies("+shared"))) + + entries.append(cmake_cache_option("RAJA_ENABLE_DESUL_ATOMICS", spec.satisfies("+desul"))) + + entries.append( + cmake_cache_option("RAJA_ENABLE_VECTORIZATION", spec.satisfies("+vectorization")) + ) + + entries.append(cmake_cache_option("RAJA_ENABLE_OPENMP_TASK", spec.satisfies("+omptask"))) + + entries.append( + cmake_cache_option("RAJA_ENABLE_TARGET_OPENMP", spec.satisfies("+omptarget")) + ) + + entries.append(cmake_cache_option("RAJA_ENABLE_SYCL", spec.satisfies("+sycl"))) + entries.append( + cmake_cache_option("RAJA_ENABLE_NV_TOOLS_EXT", spec.satisfies("+gpu-profiling +cuda")) + ) + entries.append( + cmake_cache_option("RAJA_ENABLE_ROCTX", spec.satisfies("+gpu-profiling +rocm")) + ) + + if spec.satisfies("+lowopttest"): + entries.append(cmake_cache_string("CMAKE_CXX_FLAGS_RELEASE", "-O1")) + + # C++ standard + entries.append(cmake_cache_string("BLT_CXX_STD", f"c++{self.cxx_std}")) + + entries.append( + cmake_cache_option("RAJA_ENABLE_RUNTIME_PLUGINS", spec.satisfies("+plugins")) + ) + + if spec.satisfies("+omptarget"): + entries.append( + cmake_cache_string( + "BLT_OPENMP_COMPILE_FLAGS", "-fopenmp;-fopenmp-targets=nvptx64-nvidia-cuda" + ) + ) + entries.append( + cmake_cache_string( + "BLT_OPENMP_LINK_FLAGS", "-fopenmp;-fopenmp-targets=nvptx64-nvidia-cuda" + ) + ) + + entries.append( + cmake_cache_option( + "{}ENABLE_EXAMPLES".format(option_prefix), spec.satisfies("+examples") + ) + ) + if spec.satisfies("@0.14.0:"): + entries.append( + cmake_cache_option( + "{}ENABLE_EXERCISES".format(option_prefix), spec.satisfies("+exercises") + ) + ) + else: + entries.append(cmake_cache_option("ENABLE_EXERCISES", spec.satisfies("+exercises"))) + + # TODO: Treat the workaround when building tests with spack wrapper + # For now, removing it to test CI, which builds tests outside of wrapper. + # Work around spack adding -march=ppc64le to SPACK_TARGET_ARGS which + # is used by the spack compiler wrapper. This can go away when BLT + # removes -Werror from GTest flags + # + # if self.spec.satisfies("%clang target=ppc64le:") + # or (not self.run_tests and not spec.satisfies("+tests")): + if not self.run_tests and not spec.satisfies("+tests"): + entries.append(cmake_cache_option("ENABLE_TESTS", False)) + else: + entries.append(cmake_cache_option("ENABLE_TESTS", True)) + if not spec.satisfies("+run-all-tests"): + if spec.satisfies("%clang@12.0.0:13.9.999"): + entries.append( + cmake_cache_string( + "CTEST_CUSTOM_TESTS_IGNORE", + "test-algorithm-sort-OpenMP.exe;test-algorithm-stable-sort-OpenMP.exe", + ) + ) + excluded_tests = [ + "test-algorithm-sort-Cuda.exe", + "test-algorithm-stable-sort-Cuda.exe", + "test-algorithm-sort-OpenMP.exe", + "test-algorithm-stable-sort-OpenMP.exe", + ] + if spec.satisfies("+cuda %clang@12.0.0:13.9.999"): + entries.append( + cmake_cache_string("CTEST_CUSTOM_TESTS_IGNORE", ";".join(excluded_tests)) + ) + if spec.satisfies("+cuda %xl@16.1.1.12"): + entries.append( + cmake_cache_string( + "CTEST_CUSTOM_TESTS_IGNORE", + "test-algorithm-sort-Cuda.exe;test-algorithm-stable-sort-Cuda.exe", + ) + ) + + entries.append(cmake_cache_option("RAJA_HOST_CONFIG_LOADED", True)) + + return entries + + def cmake_args(self): + return [] + + @run_after("build") + @on_package_attributes(run_tests=True) + def check_build(self): + """Run RAJA's unit test target after build when tests are enabled.""" + with working_dir(self.build_directory): + print("Running RAJA Unit Tests...") + make("test") + + examples_src_dir = "examples" + using_with_cmake_dir = join_path("examples", "using-with-cmake") + + def _rewrite_host_config(self, path): + """Replace compiler wrappers in cached install-test files.""" + kwargs = {"backup": False, "ignore_absent": True} + compiler_paths = { + "CMAKE_C_COMPILER": getattr(self.compiler, "cc", None), + "CMAKE_CXX_COMPILER": getattr(self.compiler, "cxx", None), + "CMAKE_Fortran_COMPILER": getattr(self.compiler, "fc", None), + "CMAKE_CUDA_HOST_COMPILER": getattr(self.compiler, "cxx", None), + } + + for key, value in compiler_paths.items(): + if value: + filter_file( + rf"set\({key}.*\)", f'set({key} "{value}" CACHE PATH "")', path, **kwargs + ) + + @run_after("install") + def setup_install_tests(self): + """Install and cache standalone test sources, using staged or build outputs + when available.""" + + cache_extra_test_sources(self, [self.examples_src_dir]) + + src_dir = join_path(self.stage.source_path, "test", "install", "using-with-cmake") + dst_dir = join_path(install_test_root(self), self.using_with_cmake_dir) + + if os.path.exists(src_dir): + shutil.rmtree(dst_dir, ignore_errors=True) + install_tree(src_dir, dst_dir) + self._rewrite_host_config(join_path(dst_dir, "host-config.cmake")) + + src_dir = join_path(self.build_directory, "examples", "using-with-cmake") + dst_dir = join_path(install_test_root(self), self.using_with_cmake_dir) + + if os.path.exists(src_dir): + install_tree(src_dir, dst_dir) + self._rewrite_host_config(join_path(dst_dir, "host-config.cmake")) + else: + tty.msg("Can't install host-config.cmake\n") + + def _run_common_check_install(self, test_dir): + """Verify that the using-with-cmake example can build against the installed + RAJA package.""" + + example_stage_dir = join_path(test_dir, "examples", "using-with-cmake") + with working_dir(join_path(example_stage_dir, "build"), create=True): + host_config = join_path("../", "host-config.cmake") + if not os.path.exists(host_config): + raise SkipTest(f"{os.path.abspath(host_config)} not found, cannot build example") + cmake_args = ["-C", host_config, "../"] + cmake = self.spec["cmake"].command + make_exe = which("make", required=True) + cmake(*cmake_args) + make_exe() + example = Executable("./using-with-cmake") + example() + make_exe("clean") + + @run_after("install") + @on_package_attributes(run_tests=True) + def check_install(self): + """Installation-time verification that the using-with-cmake example can build against the + installed RAJA package.""" + + src_dir = join_path(install_test_root(self)) + dst_dir = join_path(self.stage.path, "spack-test") + + if os.path.exists(src_dir): + install_tree(src_dir, dst_dir) + self._run_common_check_install(dst_dir) + else: + raise SkipTest("examples directory not found, cannot build example") + + def test_check_install(self): + """Stand-alone verification that the using-with-cmake example can build against the + installed RAJA package.""" + + self._run_common_check_install(self.test_suite.current_test_cache_dir) + + def _write_example_cmakelists(self, path, exe, source): + cmake_contents = dedent(f"""\ + cmake_minimum_required(VERSION 3.23) + project(raja_package_test LANGUAGES CXX) + + if(NOT DEFINED RAJA_DIR OR NOT EXISTS + ${{RAJA_DIR}}/lib/cmake/raja/raja-config.cmake) + message(FATAL_ERROR "Missing required 'RAJA_DIR' variable pointing to + an installed RAJA") + endif() + + find_package(RAJA REQUIRED + NO_DEFAULT_PATH + PATHS ${{RAJA_DIR}}/lib/cmake/raja) + + add_executable({exe} ../{source}) + target_link_libraries({exe} RAJA) + """) + + with open(path, "w", encoding="utf-8") as f: + f.write(cmake_contents) + + def build_and_run_example(self, exe, expected): + """Build an example from the cached test sources and verify its output.""" + + examples_dir = join_path(self.test_suite.current_test_cache_dir, self.examples_src_dir) + build_dir = join_path(examples_dir, f"build-{exe}") + with working_dir(build_dir, create=True): + cmake = self.spec["cmake"].command + make_exe = which("make", required=True) + host_config = join_path("../using-with-cmake", "host-config.cmake") + if not os.path.exists(host_config): + raise SkipTest("host-config.cmake not found, cannot build example") + self._write_example_cmakelists("CMakeLists.txt", exe, f"{exe}.cpp") + cmake_args = ["-C", host_config, "."] + cmake(*cmake_args) + make_exe() + exe_path = join_path(".", exe) + if not os.path.exists(exe_path): + raise SkipTest(f"{exe} was not built") + example = Executable(exe_path) + out = example(output=str, error=str) + check_outputs(expected, out) + make_exe("clean") + + def test_daxpy(self): + """Check daxpy tutorial""" + self.build_and_run_example("tut_daxpy", [r"daxpy", r"result -- PASS"]) + + # TODO: this test seems to hang or take a long time? + # SGS 2026-05-22: Did not see hangs/long execution times on LC systems or Redhat workstation + # clarify with Cody where this was occuring. + # def test_matrix_multiply(self): + # """check batched matrix multiple tutorial""" + # self.build_and_run_example( + # "tut_matrix-multiply", [r"matrix multiplication", r"result -- PASS"] + # ) + + def test_launch_basic(self): + """Check basic raja::launch tutorial.""" + if "+cuda" in self.spec or "+rocm" in self.spec: + self.build_and_run_example( + "tut_launch_basic", [r"Running RAJA-Teams", r"result -- PASS"] + ) + else: + raise SkipTest("CUDA or ROCm support is required to run this example") + + def test_halo_exchange(self): + """Check halo exchange tutorial.""" + self.build_and_run_example( + "tut_halo-exchange", [r"RAJA halo exchange example", r"result -- PASS"] + ) + + def test_wave_equation(self): + """Check wave equation.""" + self.build_and_run_example("wave-eqn", [r"Max Error = 2", r"Evolved solution to time"]) diff --git a/spack_repo/raja/packages/raja/tile-iterator-comparison-fix-2024.02.patch b/spack_repo/raja/packages/raja/tile-iterator-comparison-fix-2024.02.patch new file mode 100644 index 0000000000..8b849e6f82 --- /dev/null +++ b/spack_repo/raja/packages/raja/tile-iterator-comparison-fix-2024.02.patch @@ -0,0 +1,20 @@ +diff --git a/include/RAJA/pattern/kernel/Tile.hpp b/include/RAJA/pattern/kernel/Tile.hpp +index 351c263..9b1d65d 100644 +--- a/include/RAJA/pattern/kernel/Tile.hpp ++++ b/include/RAJA/pattern/kernel/Tile.hpp +@@ -169,13 +169,13 @@ struct IterableTiler { + } + + RAJA_HOST_DEVICE +- RAJA_INLINE bool operator!=(const IterableTiler &rhs) const ++ RAJA_INLINE bool operator!=(const iterator &rhs) const + { + return block_id != rhs.block_id; + } + + RAJA_HOST_DEVICE +- RAJA_INLINE bool operator<(const IterableTiler &rhs) const ++ RAJA_INLINE bool operator<(const iterator &rhs) const + { + return block_id < rhs.block_id; + } diff --git a/spack_repo/raja/repo.yaml b/spack_repo/raja/repo.yaml new file mode 100644 index 0000000000..51a9e7db91 --- /dev/null +++ b/spack_repo/raja/repo.yaml @@ -0,0 +1,3 @@ +repo: + namespace: raja + api: v2.2 From 3c0897225ce10407ba945c68e6948ba94a84be2c Mon Sep 17 00:00:00 2001 From: Brandon Echols Date: Wed, 15 Jul 2026 15:24:45 -0700 Subject: [PATCH 12/15] Empty commit for CI From 5f35874fe41735f210ffc17dbc235ea6749a659a Mon Sep 17 00:00:00 2001 From: Brandon Echols Date: Thu, 16 Jul 2026 13:51:11 -0700 Subject: [PATCH 13/15] Makes SYCL reduce more similar to CUDA and HIP backends - Fixes a segfault - More work is needed to better support SYCL's reduce and be more similar to CUDA and HIP. --- include/RAJA/policy/sycl/MemUtils_SYCL.hpp | 6 +++++- include/RAJA/policy/sycl/reduce.hpp | 8 ++++---- 2 files changed, 9 insertions(+), 5 deletions(-) diff --git a/include/RAJA/policy/sycl/MemUtils_SYCL.hpp b/include/RAJA/policy/sycl/MemUtils_SYCL.hpp index fc100449dd..beeebd78de 100644 --- a/include/RAJA/policy/sycl/MemUtils_SYCL.hpp +++ b/include/RAJA/policy/sycl/MemUtils_SYCL.hpp @@ -54,7 +54,7 @@ struct syclInfo { sycl_dim_t gridDim {0}; sycl_dim_t blockDim {0}; - ::sycl::queue qu = ::sycl::queue(); + ::sycl::queue qu = ::sycl::queue(); bool setup_reducers = false; }; @@ -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 { diff --git a/include/RAJA/policy/sycl/reduce.hpp b/include/RAJA/policy/sycl/reduce.hpp index 26528af129..8429c4381e 100644 --- a/include/RAJA/policy/sycl/reduce.hpp +++ b/include/RAJA/policy/sycl/reduce.hpp @@ -120,7 +120,7 @@ struct Reduce_Data Reduce_Data(T initValue, T identityValue, Offload_Info& info) : value(initValue) { - ::sycl::queue& q = ::camp::resources::Sycl::get_default().get_queue(); + ::sycl::queue q = currentResourceQueue(); device = reinterpret_cast( @@ -153,7 +153,7 @@ struct Reduce_Data //! transfers from the host to the device -- exit() is called upon failure RAJA_INLINE void hostToDevice(Offload_Info& RAJA_UNUSED_ARG(info)) { - ::sycl::queue& q = ::camp::resources::Sycl::get_default().get_queue(); + ::sycl::queue q = currentResourceQueue(); // precondition: host and device are valid pointers auto e = @@ -166,7 +166,7 @@ struct Reduce_Data //! transfers from the device to the host -- exit() is called upon failure RAJA_INLINE void deviceToHost(Offload_Info& RAJA_UNUSED_ARG(info)) { - ::sycl::queue& q = ::camp::resources::Sycl::get_default().get_queue(); + ::sycl::queue q = currentResourceQueue(); // precondition: host and device are valid pointers auto e = @@ -179,7 +179,7 @@ struct Reduce_Data //! frees all data from the offload information passed RAJA_INLINE void cleanup(Offload_Info& RAJA_UNUSED_ARG(info)) { - ::sycl::queue& q = ::camp::resources::Sycl::get_default().get_queue(); + ::sycl::queue q = currentResourceQueue(); if (device) { From bd8b8a66e3d73384ae4af3adb97dcd5c94c0881c Mon Sep 17 00:00:00 2001 From: Brandon Echols Date: Thu, 16 Jul 2026 15:37:52 -0700 Subject: [PATCH 14/15] Returns queue as reference and formats code --- include/RAJA/policy/sycl/MemUtils_SYCL.hpp | 4 ++-- include/RAJA/policy/sycl/reduce.hpp | 8 ++++---- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/include/RAJA/policy/sycl/MemUtils_SYCL.hpp b/include/RAJA/policy/sycl/MemUtils_SYCL.hpp index beeebd78de..22587251e8 100644 --- a/include/RAJA/policy/sycl/MemUtils_SYCL.hpp +++ b/include/RAJA/policy/sycl/MemUtils_SYCL.hpp @@ -54,7 +54,7 @@ struct syclInfo { sycl_dim_t gridDim {0}; sycl_dim_t blockDim {0}; - ::sycl::queue qu = ::sycl::queue(); + ::sycl::queue qu = ::sycl::queue(); bool setup_reducers = false; }; @@ -68,7 +68,7 @@ extern std::unordered_map<::sycl::queue, bool> g_queue_info_map; //! get queue for current launch RAJA_INLINE -::sycl::queue currentResourceQueue() { return detail::tl_status.qu; } +::sycl::queue& currentResourceQueue() { return detail::tl_status.qu; } //! Allocator for pinned memory for use in basic_mempool struct PinnedAllocator diff --git a/include/RAJA/policy/sycl/reduce.hpp b/include/RAJA/policy/sycl/reduce.hpp index 8429c4381e..0e7fe6d3f8 100644 --- a/include/RAJA/policy/sycl/reduce.hpp +++ b/include/RAJA/policy/sycl/reduce.hpp @@ -120,7 +120,7 @@ struct Reduce_Data Reduce_Data(T initValue, T identityValue, Offload_Info& info) : value(initValue) { - ::sycl::queue q = currentResourceQueue(); + ::sycl::queue& q = currentResourceQueue(); device = reinterpret_cast( @@ -153,7 +153,7 @@ struct Reduce_Data //! transfers from the host to the device -- exit() is called upon failure RAJA_INLINE void hostToDevice(Offload_Info& RAJA_UNUSED_ARG(info)) { - ::sycl::queue q = currentResourceQueue(); + ::sycl::queue& q = currentResourceQueue(); // precondition: host and device are valid pointers auto e = @@ -166,7 +166,7 @@ struct Reduce_Data //! transfers from the device to the host -- exit() is called upon failure RAJA_INLINE void deviceToHost(Offload_Info& RAJA_UNUSED_ARG(info)) { - ::sycl::queue q = currentResourceQueue(); + ::sycl::queue& q = currentResourceQueue(); // precondition: host and device are valid pointers auto e = @@ -179,7 +179,7 @@ struct Reduce_Data //! frees all data from the offload information passed RAJA_INLINE void cleanup(Offload_Info& RAJA_UNUSED_ARG(info)) { - ::sycl::queue q = currentResourceQueue(); + ::sycl::queue& q = currentResourceQueue(); if (device) { From 637cd1b6352f0132a2cf40a84d1a9f36c99e974f Mon Sep 17 00:00:00 2001 From: Rich Hornung Date: Mon, 20 Jul 2026 12:52:06 -0700 Subject: [PATCH 15/15] Try a stable version of cmake in windows build --- .github/workflows/build.yml | 1 + 1 file changed, 1 insertion(+) diff --git a/.github/workflows/build.yml b/.github/workflows/build.yml index 86ae718008..804b3f1f60 100644 --- a/.github/workflows/build.yml +++ b/.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