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 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/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/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/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; } 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 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 diff --git a/include/RAJA/policy/sycl/MemUtils_SYCL.hpp b/include/RAJA/policy/sycl/MemUtils_SYCL.hpp index 13c941779f..22587251e8 100644 --- a/include/RAJA/policy/sycl/MemUtils_SYCL.hpp +++ b/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,8 +78,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 +87,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 +101,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 +110,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,9 +125,9 @@ 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; } @@ -131,8 +135,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..0e7fe6d3f8 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 = currentResourceQueue(); 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) { @@ -147,21 +147,18 @@ 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)) { - ::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 = currentResourceQueue(); // 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 +166,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 = currentResourceQueue(); // 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 +179,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 = currentResourceQueue(); 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; } 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 \ 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 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); 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