diff --git a/.github/workflows/build.yml b/.github/workflows/build.yml index 86ae718008..11dcea117b 100644 --- a/.github/workflows/build.yml +++ b/.github/workflows/build.yml @@ -68,6 +68,61 @@ jobs: ## - run: | ## dir -r D:\a\RAJA\RAJA\build\test ## ==================================== +## Run tests action + - uses: threeal/ctest-action@v1.1.0 + with: + build-config: Debug + + build_windows_cuda: + strategy: + matrix: + shared: + - args: + BUILD_SHARED_LIBS=Off + - args: + BUILD_SHARED_LIBS=On + CMAKE_WINDOWS_EXPORT_ALL_SYMBOLS=On + + # Running on Visual Studio 2022 for better CUDA 12 compatibility + # windows-2025 would be locked to Visual Studio 2026 + # windows-latest can be changed by GitHub + runs-on: windows-2022 + + steps: + - uses: actions/checkout@v2 + with: + submodules: recursive + + - uses: Jimver/cuda-toolkit@v0.2.35 + with: + cuda: '12.9.1' + sub-packages: '["nvcc", "visual_studio_integration", "cudart", "nvrtc", "thrust"]' + + - name: Check CUDA + shell: pwsh + run: | + nvcc --version + echo "CUDA_PATH=$env:CUDA_PATH" + + - uses: threeal/cmake-action@v1.3.0 + with: + build-dir: build + options: + ENABLE_WARNINGS_AS_ERRORS=Off + BLT_CXX_STD="c++20" + CMAKE_CUDA_FLAGS="-Xcompiler=/TP" + CMAKE_CXX_STANDARD=20 + CMAKE_BUILD_TYPE=Release + CMAKE_CUDA_ARCHITECTURES=75 + ENABLE_CUDA=On + ${{ matrix.shared.args }} + run-build: true +# build-args: '--parallel 16' +## ==================================== +## Print the contents of the test directory in the build space (debugging) +## - run: | +## dir -r D:\a\RAJA\RAJA\build\test +## ==================================== ## Run tests action - uses: threeal/ctest-action@v1.1.0 with: diff --git a/.gitlab/jobs/matrix.yml b/.gitlab/jobs/matrix.yml index 73a8eef9e1..96c1c88e0c 100644 --- a/.gitlab/jobs/matrix.yml +++ b/.gitlab/jobs/matrix.yml @@ -36,3 +36,8 @@ clang_18_1_8_cuda_12_6_0: variables: SPEC: " ~shared +cuda cuda_arch=90 +tests %llvm@=18.1.8 ^cuda@12.6.0+allow-unsupported-compilers" extends: .job_on_matrix + +clang_18_1_8_cuda_13_1_1: + variables: + SPEC: " ~shared +cuda cuda_arch=90 +tests %llvm@=18.1.8 ^cuda@13.1.1+allow-unsupported-compilers" + extends: .job_on_matrix diff --git a/CMakeLists.txt b/CMakeLists.txt index df84df2d7c..f96a11a023 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -343,6 +343,12 @@ endif() if(RAJA_ENABLE_TESTS) add_subdirectory(test) +else() + # For Spack stand-alone tests need host config information for building simple + # examples even if full tests are not done. + configure_file( + test/install/using-with-cmake/host-config.cmake.in + ${PROJECT_BINARY_DIR}/examples/using-with-cmake/host-config.cmake) endif() if(RAJA_ENABLE_REPRODUCERS) diff --git a/blt b/blt index e783e30f28..f08f95e0f4 160000 --- a/blt +++ b/blt @@ -1 +1 @@ -Subproject commit e783e30f2823ee1a208f7f90741b41c1f5a08063 +Subproject commit f08f95e0f49b1d53db8be2dd577d1cdbfd9e9f90 diff --git a/cmake/thirdparty/FindCUB.cmake b/cmake/thirdparty/FindCUB.cmake index 866157ff85..53bf3cd3b9 100644 --- a/cmake/thirdparty/FindCUB.cmake +++ b/cmake/thirdparty/FindCUB.cmake @@ -14,7 +14,8 @@ find_path(CUB_INCLUDE_DIRS HINTS ${CUB_DIR}/ ${CUB_DIR}/include - ${CUDA_TOOLKIT_ROOT_DIR}/include) + ${CUDA_TOOLKIT_ROOT_DIR}/include + ${CUDA_TOOLKIT_ROOT_DIR}/include/cccl) find_package_handle_standard_args( CUB diff --git a/include/RAJA/policy/cuda/MemUtils_CUDA.hpp b/include/RAJA/policy/cuda/MemUtils_CUDA.hpp index 6b51d04249..0812b24680 100644 --- a/include/RAJA/policy/cuda/MemUtils_CUDA.hpp +++ b/include/RAJA/policy/cuda/MemUtils_CUDA.hpp @@ -146,10 +146,26 @@ struct DevicePinnedAllocator void* ptr; CAMP_CUDA_API_INVOKE_AND_CHECK(cudaMallocManaged, &ptr, nbytes, cudaMemAttachGlobal); +#if CUDART_VERSION >= 13000 + cudaMemLocation devLoc {}; + devLoc.type = cudaMemLocationTypeDevice; + devLoc.id = device; + + CAMP_CUDA_API_INVOKE_AND_CHECK(cudaMemAdvise, ptr, nbytes, + cudaMemAdviseSetPreferredLocation, devLoc); + + cudaMemLocation cpuLoc {}; + cpuLoc.type = cudaMemLocationTypeHost; + cpuLoc.id = 0; + + CAMP_CUDA_API_INVOKE_AND_CHECK(cudaMemAdvise, ptr, nbytes, + cudaMemAdviseSetAccessedBy, cpuLoc); +#else CAMP_CUDA_API_INVOKE_AND_CHECK(cudaMemAdvise, ptr, nbytes, cudaMemAdviseSetPreferredLocation, device); CAMP_CUDA_API_INVOKE_AND_CHECK(cudaMemAdvise, ptr, nbytes, cudaMemAdviseSetAccessedBy, cudaCpuDeviceId); +#endif return ptr; } diff --git a/scripts/radiuss-spack-configs b/scripts/radiuss-spack-configs index 3a7d108251..5a73b91c06 160000 --- a/scripts/radiuss-spack-configs +++ b/scripts/radiuss-spack-configs @@ -1 +1 @@ -Subproject commit 3a7d108251fa3d7c014a4b74bf22d8cf181411ee +Subproject commit 5a73b91c06724ff089c13db1a3d197b5ea9d5c0c diff --git a/test/unit/resource/tests/test-resource-AsyncTime.hpp b/test/unit/resource/tests/test-resource-AsyncTime.hpp index d3cede75e7..47c9462115 100644 --- a/test/unit/resource/tests/test-resource-AsyncTime.hpp +++ b/test/unit/resource/tests/test-resource-AsyncTime.hpp @@ -32,22 +32,31 @@ gpu_time_wait_for(float time, float clockrate) { int get_clockrate() { int cuda_device = 0; - cudaDeviceProp deviceProp; CAMP_CUDA_API_INVOKE_AND_CHECK(cudaGetDevice, &cuda_device); - CAMP_CUDA_API_INVOKE_AND_CHECK(cudaGetDeviceProperties, &deviceProp, cuda_device); - if ((deviceProp.concurrentKernels == 0)) + + int concurrentKernels = 0; + CAMP_CUDA_API_INVOKE_AND_CHECK(cudaDeviceGetAttribute, + &concurrentKernels, + cudaDevAttrConcurrentKernels, + cuda_device); + + if (concurrentKernels == 0) { printf("> GPU does not support concurrent kernel execution\n"); printf(" CUDA kernel runs will be serialized\n"); return -1; } - //printf("> Detected Compute SM %d.%d hardware with %d multi-processors\n", - // deviceProp.major, deviceProp.minor, deviceProp.multiProcessorCount); + + int clockRate = 0; + CAMP_CUDA_API_INVOKE_AND_CHECK(cudaDeviceGetAttribute, + &clockRate, + cudaDevAttrClockRate, + cuda_device); #if defined(__arm__) || defined(__aarch64__) - return deviceProp.clockRate/1000; + return clockRate / 1000; #else - return deviceProp.clockRate; + return clockRate; #endif } diff --git a/tmp-windows-cuda-notes/README.md b/tmp-windows-cuda-notes/README.md new file mode 100644 index 0000000000..1b2f29b561 --- /dev/null +++ b/tmp-windows-cuda-notes/README.md @@ -0,0 +1,192 @@ +# RAJA Build with CUDA on Windows + +## Window CUDA CI build + +A Windows CUDA CI job was added. Not working since the RAJA compile is not working. + +See ../.github/workflows/build.yml + +Should adjust CUDA and MSVCC versions when a working compile. + +Do we want to test multiple CUDA MSVCC versions? + +## MS toolchain + +### Install Visual Studio + +2026 https://visualstudio.microsoft.com/downloads/ +2022 https://aka.ms/vs/17/release/vs\_community.exe + +Steve : tried both initially to see if one would show fewer errors. The errors seem to be mostly the same for both so used 2026. + +NOTE: MS seems to remove references to old VS community versions, found above 2022 link on reddit not sure how long links remain for the old versions. +Paid subscribers can get access to the old VS versions. + +### Install Cmake + +https://cmake.org/download/ + +### Install Git + +https://git-scm.com/install/windows + +### Install CUDA + +https://developer.nvidia.com/cuda-downloads + +Steve: Tried both CUDA 12 and CUDA 13. Both had issues so went with CUDA 13 since we have added support for it. + +## Git stuff + +All generic stuff. Added a key for GitHub. + +``` +ssh-keygen -t ed25519 -C "smith84@llnl.gov" +type %userprofile%\\.ssh\\id\_ed25519.pub | clip + +git clone --recursive git@github.com:llnl/RAJA.git +``` + +## Failed attempts to build. + +Steve : Tried combinations of MSVCC 2022 and 2026 with Cuda 12 and 13 to see if some combination might work....needed to start deeper patching. + +## Steve patching + +### BLT + +Apply the blt.patch to fix argument passing to MSVCC. nvcc does not pass "/arg" to MSVCC by default which causes nvcc to interpret the arguments as directories/files. + +Kenny has similar patch for blt in AXOM branch. + +``` + +cd blt + +git apply ..\\tmp-windows-cuda-notes\\blt.patch + +``` + +### Camp + +MSVCC has issue parsing variadic macros. Uses non-standard CPP parser. Did come up with patch for CAMP but Nvidia CUDA 13 recommends using the "-Xcompiler=/Zc:preprocessor" flag which uses more standard conforming CPP. Seemed more reasonable approach. + +Camp patch is in camp-variadic.patch. + +## RAJA config/build + +CUDA 13 recommends the preprocessor flag and fix a CAMP issue. MSVCC uses some non-standard conforming parsing if this is not done. + +"/TP" treats .C files as C++. Was fixing some issues with CUDA 12 toolkit during initial attempts but should double check this requirement. + +``` +cmake -B cudabuild RAJA -DENABLE\_CUDA=ON -DCMAKE\_CUDA\_FLAGS=" --keep -Xcompiler=/TP -Xcompiler=/Zc:preprocessor" -DCMAKE\_CUDA\_ARCHITECTURES=75 + +``` + +``` +cmake --build cudabuild --target atomic-histogram\_solution > zz.build.txt 2>\&1 +``` + +Steve: reached a point with a template issue, which above build exposes. + +``` + atomic-histogram\_solution.cpp + atomic-histogram\_solution.cudafe1.cpp +C:\\Users\\smith84\\projects\\raja\\clean\\RAJA\\include\\RAJA/index/IndexSet.hpp(754): warning C4244: 'return': conversion from '\_\_int64' to 'int', possible + loss of data \[C:\\Users\\smith84\\projects\\raja\\clean\\cudabuild\\exercises\\atomic-histogram\_solution.vcxproj] +C:\\Users\\smith84\\projects\\raja\\clean\\RAJA\\include\\RAJA/index/IndexSet.hpp(759): warning C4244: 'return': conversion from 'const \_\_int64' to 'int', po +ssible loss of data \[C:\\Users\\smith84\\projects\\raja\\clean\\cudabuild\\exercises\\atomic-histogram\_solution.vcxproj] +C:\\Users\\smith84\\projects\\raja\\clean\\cudabuild\\exercises\\atomic-h.9E0645A7\\x64\\Debug\\atomic-histogram\_solution.cudafe1.stub.c(37): error C2912: expli +cit specialization 'void RAJA::policy::cuda::impl::\_\_wrapper\_\_device\_stub\_forallp\_cuda\_kernel,RAJA::CudaDefaultConcretizer,1,false>,1,RAJA::Iterators::numeric\_iterator,\_\_nv\_dl\_wrapper\_t<\_\_nv\_dl\_tag,int \*,int \*>,int,RAJA::expt::ForallParamPack<>,RAJ +A::iteration\_mapping::Direct,RAJA::cuda::IndexGlobal,256>(const \_ZN4camp5decayIZ4mainEUliE0\_EE \&,const \_ZN4camp5decayIN4RAJ +A9Iterators16numeric\_iteratorIiiPiEEEE \&,const \_ZN4camp5decayIiEE \&,\_ZN4camp5decayIN4RAJA4expt15ForallParamPackIJEEEEE \&)' is not a specialization of + a function template \[C:\\Users\\smith84\\projects\\raja\\clean\\cudabuild\\exercises\\atomic-histogram\_solution.vcxproj] + with + \[ + StorageT=int, + DiffT=int, + Type=int + ] +C:\\Program Files\\Microsoft Visual Studio\\18\\Community\\MSBuild\\Microsoft\\VC\\v180\\BuildCustomizations\\CUDA 13.3.targets(804,9): error MSB3721: The comm +and ""C:\\Program Files\\NVIDIA GPU Computing Toolkit\\CUDA\\v13.3\\bin\\nvcc.exe" --use-local-env -ccbin "C:\\Program Files\\Microsoft Visual Studio\\18\\Com +munity\\VC\\Tools\\MSVC\\14.51.36231\\bin\\HostX64\\x64" -x cu -IC:\\Users\\smith84\\projects\\raja\\clean\\RAJA\\include -IC:\\Users\\smith84\\projects\\raja\\clean\\ +cudabuild\\include -IC:\\Users\\smith84\\projects\\raja\\clean\\RAJA\\tpl\\camp\\include -IC:\\Users\\smith84\\projects\\raja\\clean\\cudabuild\\tpl\\camp\\include -I"C +:\\Program Files\\NVIDIA GPU Computing Toolkit\\CUDA\\v13.3\\include" -I"C:\\Program Files\\NVIDIA GPU Computing Toolkit\\CUDA\\v13.3\\include\\cccl" -I"C:\\Prog +ram Files\\NVIDIA GPU Computing Toolkit\\CUDA\\v13.3\\include" -G --keep --keep-dir atomic-h.9E0645A7\\x64\\Debug -maxrregcount=0 --machine 64 --com +pile -cudart static -forward-unknown-to-host-compiler -restrict --expt-extended-lambda --expt-relaxed-constexpr -Xcudafe --display\_error\_number -O0 - +std=c++20 --generate-code=arch=compute\_75,code=\[compute\_75,sm\_75] -Xcompiler="/TP /Zc:preprocessor" -g -DRAJA\_WIN\_STATIC\_BUILD -DCAMP\_WIN\_STATIC\_BUI +LD -D"CMAKE\_INTDIR=\\"Debug\\"" -D\_MBCS -DWIN32 -D\_WINDOWS -DRAJA\_WIN\_STATIC\_BUILD -DCAMP\_WIN\_STATIC\_BUILD -D"CMAKE\_INTDIR=\\"Debug\\"" -Xcompiler "/EHsc + /W3 /nologo /FS /Zi /RTC1 /MTd /GR" -Xcompiler "/Fdatomic-histogram\_solution.dir\\Debug\\vc145.pdb" -o atomic-histogram\_solution.dir\\Debug\\atomic-his +togram\_solution.obj "C:\\Users\\smith84\\projects\\raja\\clean\\RAJA\\exercises\\atomic-histogram\_solution.cpp"" exited with code 2. \[C:\\Users\\smith84\\projec +ts\\raja\\clean\\cudabuild\\exercises\\atomic-histogram\_solution.vcxproj] +``` + +## Kenny Weiss work on AXOM + +Steve : Learned of Kenny branch of AXOM. Kenny took Steve's CUDA 13 branch and made patches with extra fixes for Windows. Steve was starting to parse through Kenny patching. The template issue above might be have been similar to what is patched in "kweiss-windows-cuda-msvc.patch" pulled from Kenny's repository. Current RAJA has modified the templates so this needs some work. + +Steve goes on vacation :( + +Kenny/AI found some other changes, note Kenny was using VS 2022 and Cuda 13. + +## Kenny Notes: + +Hi Steve, + +I do! + +(CC-ing Rich and Brian Han since they’re also likely interested in the details) + +I’m still working through some of the details, so I haven’t created a PR yet, but the branch is here: https://github.com/llnl/axom/compare/develop...feature/kweiss/windows-gpu + +I used codex w/ gpt@5.5 to work through a bunch of the issues. The compilation times on my laptop (6 core i7; I forget the GPU, but it was CUDA\_ARCHITECTURE 75) can be really long, so each round took a few hours, but it eventually got there. + +Axom uses vcpkg (https://github.com/microsoft/vcpkg) through uberenv ( https://github.com/llnl/uberenv/ | https://uberenv.readthedocs.io/en/latest/) + +to manage our TPLs on windows. We manage some local ports of raja, umpire, mfem, conduit, … inside of axom’s ./scripts/vcpkg\_ports directory. + +(ports are equivalent to spack package recipes) + +To get this working w/ uberenv, I had to add some slight edits: https://github.com/llnl/uberenv/pull/160 -- which should be merged soon + +I also have a few changes to BLT, but I haven’t yet looked closely enough to determine which are strictly necessary, and haven’t posted a PR/branch to BLT’s github yet + +My branch has a bunch of updates to our vcpkg ports to get raja, umpire and camp to build on msvc+cuda. + +Since my build uses cuda@13, I needed some patches from your raja and camp PRs related to cuda@13. + +To set up my build environment on Windows, I had to install: + +the cuda toolkit ( https://developer.nvidia.com/cuda-downloads -- I am using cuda@13) + +MS visual studio community ( https://visualstudio.microsoft.com/vs/community/ -- + +I am using visual studio community 2022. Presumably it’d work w/ visual studio community 2026 ). + +If you need MPI, you also need to install MS MPI -- https://learn.microsoft.com/en-us/message-passing-interface/microsoft-mpi + +Some gotchas in my msvc+cuda porting: + +cuda@13 has some changes (but you’re clearly aware of those) + +MSVC does not support extended lambdas inside of if constexpr -- + +See item 11 in Section 5.3.8.4 -- https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-support.html#extended-lambda-restrictions + +Microsoft defines some regular words as macros, e.g. min, max, small, large + +I hope this helps and would be happy to discuss this further/answer questions, + +Kenny + +\-- + +Kenneth Weiss + +Numerics, Modular Applications and Performance Group Leader + +ASQ Division, Computing Directorate, LLNL + diff --git a/tmp-windows-cuda-notes/blt.patch b/tmp-windows-cuda-notes/blt.patch new file mode 100644 index 0000000000..048ec45f13 --- /dev/null +++ b/tmp-windows-cuda-notes/blt.patch @@ -0,0 +1,28 @@ +diff --git a/thirdparty_builtin/CMakeLists.txt b/thirdparty_builtin/CMakeLists.txt +index 74a4f8d..56c59a3 100644 +--- a/thirdparty_builtin/CMakeLists.txt ++++ b/thirdparty_builtin/CMakeLists.txt +@@ -111,10 +111,19 @@ if(ENABLE_TESTS) + # + # For more info see: https://github.com/LLNL/blt/issues/79 + # +- +- blt_append_custom_compiler_flag( FLAGS_VAR gtest_extra_flags +- DEFAULT " " +- MSVC "/wd4251" "/wd4275" ) ++ if(MSVC AND ENABLE_CUDA) ++ # For CUDA nvcc compiler pass arg to MSVCC with -Xcompiler ++ set(gtest_extra_flags ++ "$<$>:/wd4251>" ++ "$<$>:/wd4275>" ++ "$<$:-Xcompiler=/wd4251>" ++ "$<$:-Xcompiler=/wd4275>" ++ ) ++ else() ++ blt_append_custom_compiler_flag( FLAGS_VAR gtest_extra_flags ++ DEFAULT " " ++ MSVC "/wd4251" "/wd4275" ) ++ endif() + + set (INSTALL_GMOCK OFF CACHE BOOL "") + set (INSTALL_GTEST OFF CACHE BOOL "") diff --git a/tmp-windows-cuda-notes/camp-variadic.patch b/tmp-windows-cuda-notes/camp-variadic.patch new file mode 100644 index 0000000000..98a74e10c1 --- /dev/null +++ b/tmp-windows-cuda-notes/camp-variadic.patch @@ -0,0 +1,24 @@ +diff --git a/include/camp/defines.hpp b/include/camp/defines.hpp +index 2d80bc3..a68f200 100644 +--- a/include/camp/defines.hpp ++++ b/include/camp/defines.hpp +@@ -219,15 +219,15 @@ cudaAssert(cudaError_t code, const char *call, const char *file, int line); + __LINE__); \ + } + /// +-#define CAMP_CUDA_API_INVOKE_AND_CHECK_RETURN(...) \ ++#define CAMP_CUDA_API_INVOKE_AND_CHECK_RETURN(func, ...) \ + [&]() { \ +- CAMP_CUDA_API_INVOKE_AND_CHECK_IMPLEMENTATION(__VA_ARGS__) \ ++ CAMP_CUDA_API_INVOKE_AND_CHECK_IMPLEMENTATION(func, __VA_ARGS__) \ + return code_56792578; \ + }() + /// +-#define CAMP_CUDA_API_INVOKE_AND_CHECK(...) \ ++#define CAMP_CUDA_API_INVOKE_AND_CHECK(func, ...) \ + do { \ +- CAMP_CUDA_API_INVOKE_AND_CHECK_IMPLEMENTATION(__VA_ARGS__) \ ++ CAMP_CUDA_API_INVOKE_AND_CHECK_IMPLEMENTATION(func, __VA_ARGS__) \ + } while (false) + + #endif // #ifdef CAMP_ENABLE_CUDA diff --git a/tmp-windows-cuda-notes/kweiss-windows-cuda-msvc.patch b/tmp-windows-cuda-notes/kweiss-windows-cuda-msvc.patch new file mode 100644 index 0000000000..5c76ba63c3 --- /dev/null +++ b/tmp-windows-cuda-notes/kweiss-windows-cuda-msvc.patch @@ -0,0 +1,206 @@ +diff --git a/include/RAJA/config.hpp.in b/include/RAJA/config.hpp.in +index 3ec07c5..4a2b6c1 100644 +--- a/include/RAJA/config.hpp.in ++++ b/include/RAJA/config.hpp.in +@@ -291,7 +291,7 @@ namespace RAJA { + #define RAJA_CUDA_DIM_T dim3 + #endif + +-#if ( (__CUDACC_VER_MAJOR__ > 11) || (__CUDACC_VER_MAJOR__ == 11 && __CUDACC_VER_MINOR__ >= 7) ) && !defined(RAJA_ENABLE_CLANG_CUDA) ++#if ( (__CUDACC_VER_MAJOR__ > 11) || (__CUDACC_VER_MAJOR__ == 11 && __CUDACC_VER_MINOR__ >= 7) ) && !defined(RAJA_ENABLE_CLANG_CUDA) && !defined(_MSC_VER) + // only supported by nvcc on sm_70 and above + #define RAJA_CUDA_GRID_CONSTANT __grid_constant__ + #else +diff --git a/cmake/SetupCompilers.cmake b/cmake/SetupCompilers.cmake +index 20c1d8f..de36e39 100644 +--- a/cmake/SetupCompilers.cmake ++++ b/cmake/SetupCompilers.cmake +@@ -27,7 +27,7 @@ endif() + + set(RAJA_COMPILER "RAJA_COMPILER_${CMAKE_CXX_COMPILER_ID}") + +-if ( MSVC ) ++if ( MSVC AND NOT DEFINED VCPKG_TARGET_TRIPLET ) + if (NOT BUILD_SHARED_LIBS) + foreach(flag_var + CMAKE_CXX_FLAGS CMAKE_CXX_FLAGS_DEBUG CMAKE_CXX_FLAGS_RELEASE +diff --git a/include/RAJA/policy/cuda/forall.hpp b/include/RAJA/policy/cuda/forall.hpp +index c1632f7..b525a00 100644 +--- a/include/RAJA/policy/cuda/forall.hpp ++++ b/include/RAJA/policy/cuda/forall.hpp +@@ -501,6 +501,114 @@ __global__ void forallp_cuda_kernel( + RAJA::expt::ParamMultiplexer::parampack_combine(EXEC_POL {}, f_params); + } + ++#if defined(_MSC_VER) ++template ++__launch_bounds__(IterationGetter::block_size, BlocksPerSM) __global__ ++ void forallp_cuda_kernel_direct_fixed( ++ const RAJA_CUDA_GRID_CONSTANT LOOP_BODY loop_body, ++ const RAJA_CUDA_GRID_CONSTANT Iterator idx, ++ const RAJA_CUDA_GRID_CONSTANT IndexType length, ++ ForallParam f_params) ++{ ++ using RAJA::internal::thread_privatize; ++ auto privatizer = thread_privatize(loop_body); ++ auto& body = privatizer.get_priv(); ++ auto ii = IterationGetter::template index(); ++ ++ if (ii < length) ++ { ++ RAJA::expt::invoke_body(f_params, body, idx[ii]); ++ } ++ ++ RAJA::expt::ParamMultiplexer::parampack_combine(EXEC_POL {}, f_params); ++} ++ ++template ++__global__ void forallp_cuda_kernel_direct_dynamic( ++ const RAJA_CUDA_GRID_CONSTANT LOOP_BODY loop_body, ++ const RAJA_CUDA_GRID_CONSTANT Iterator idx, ++ const RAJA_CUDA_GRID_CONSTANT IndexType length, ++ ForallParam f_params) ++{ ++ using RAJA::internal::thread_privatize; ++ auto privatizer = thread_privatize(loop_body); ++ auto& body = privatizer.get_priv(); ++ auto ii = IterationGetter::template index(); ++ ++ if (ii < length) ++ { ++ RAJA::expt::invoke_body(f_params, body, idx[ii]); ++ } ++ ++ RAJA::expt::ParamMultiplexer::parampack_combine(EXEC_POL {}, f_params); ++} ++ ++template ++__launch_bounds__(IterationGetter::block_size, BlocksPerSM) __global__ ++ void forallp_cuda_kernel_strided_fixed( ++ const RAJA_CUDA_GRID_CONSTANT LOOP_BODY loop_body, ++ const RAJA_CUDA_GRID_CONSTANT Iterator idx, ++ const RAJA_CUDA_GRID_CONSTANT IndexType length, ++ ForallParam f_params) ++{ ++ using RAJA::internal::thread_privatize; ++ auto privatizer = thread_privatize(loop_body); ++ auto& body = privatizer.get_priv(); ++ ++ for (auto ii = IterationGetter::template index(); ii < length; ++ ii += IterationGetter::template size()) ++ { ++ RAJA::expt::invoke_body(f_params, body, idx[ii]); ++ } ++ ++ RAJA::expt::ParamMultiplexer::parampack_combine(EXEC_POL {}, f_params); ++} ++ ++template ++__global__ void forallp_cuda_kernel_strided_dynamic( ++ const RAJA_CUDA_GRID_CONSTANT LOOP_BODY loop_body, ++ const RAJA_CUDA_GRID_CONSTANT Iterator idx, ++ const RAJA_CUDA_GRID_CONSTANT IndexType length, ++ ForallParam f_params) ++{ ++ using RAJA::internal::thread_privatize; ++ auto privatizer = thread_privatize(loop_body); ++ auto& body = privatizer.get_priv(); ++ ++ for (auto ii = IterationGetter::template index(); ii < length; ++ ii += IterationGetter::template size()) ++ { ++ RAJA::expt::invoke_body(f_params, body, idx[ii]); ++ } ++ ++ RAJA::expt::ParamMultiplexer::parampack_combine(EXEC_POL {}, f_params); ++} ++#endif ++ + } // namespace impl + + // +@@ -555,9 +663,60 @@ forall_impl(resources::Cuda cuda_res, + if (len > 0) + { + ++#if defined(_MSC_VER) ++ const void* func = nullptr; ++ if constexpr (std::is_base_of::value && ++ (IterationGetter::block_size > 0)) ++ { ++ func = reinterpret_cast( ++ &impl::forallp_cuda_kernel_direct_fixed>); ++ } ++ else if constexpr (std::is_base_of::value) ++ { ++ func = reinterpret_cast( ++ &impl::forallp_cuda_kernel_direct_dynamic>); ++ } ++ else if constexpr (std::is_base_of::value && ++ std::is_base_of::value && ++ (IterationGetter::block_size > 0)) ++ { ++ func = reinterpret_cast( ++ &impl::forallp_cuda_kernel_strided_fixed>); ++ } ++ else ++ { ++ func = reinterpret_cast( ++ &impl::forallp_cuda_kernel_strided_dynamic>); ++ } ++#else + auto func = reinterpret_cast( + &impl::forallp_cuda_kernel>); ++#endif + + // + // Setup shared memory buffers diff --git a/tpl/camp b/tpl/camp index e75ab64c02..81f341009d 160000 --- a/tpl/camp +++ b/tpl/camp @@ -1 +1 @@ -Subproject commit e75ab64c029aa27c80593715cb2a3ccad7453c8c +Subproject commit 81f341009d8080c18bb335d62e9975464948fc30