Basic kokkos #206

ajpowelsnl · 2022-01-13T19:01:30Z

This PR represents a split from PR #188. Here, we provide Kokkos basic variants. Also, we made minimal extensions to the build system for Kokkos integration. Moving forward, we will have a PR for Kokkos variants of each kernel group.

crtrott · 2022-01-24T20:03:31Z

@davidbeckingsale hey David can you take a look at this?

rhornung67 · 2022-01-25T21:12:00Z

CMakeLists.txt

+
+# The statement below is required for Kokkos compilation.
+if(ENABLE_KOKKOS)
+  include_directories(SYSTEM ${CMAKE_CURRENT_SOURCE_DIR}/tpl/RAJA/include/)


is this for basic type parameterization that the perf suite inherits from RAJA?

We removed the include_directories(SYSTEM ${CMAKE_CURRENT_SOURCE_DIR}/tpl/RAJA/include/) line from this bloc. This line was likely inserted as a hack in the early prototyping stages.

src/basic-kokkos/DAXPY-Kokkos.cpp

src/basic-kokkos/IF_QUAD-Kokkos.cpp

src/basic-kokkos/INIT3-Kokkos.cpp

src/basic-kokkos/MULADDSUB-Kokkos.cpp

src/basic-kokkos/REDUCE3_INT-Kokkos.cpp

src/basic-kokkos/TRAP_INT-Kokkos.cpp

src/common/RAJAPerfSuite.hpp

CMakeLists.txt

src/basic-kokkos/CMakeLists.txt

crtrott · 2022-02-08T00:25:04Z

Amy is gonna revert the formatting changes on the couple original files where she applied clang-format accidentally. Other than that I think this is good now.

ajpowelsnl · 2022-02-09T16:44:10Z

Good morning @davidbeckingsale -- The commits this morning (for DAXPY_ATOMIC.hpp, DAXPY_ATOMIC.cpp) suggest significant formatting changes, relative to the version of the files we're comparing to in this PR. This difference is, however, a bit misleading, because I actually took the version of the these files from your most recent develop branch. Just give a shout if you have any questions, and many thanks for your help.

crtrott · 2022-02-21T21:47:22Z

@davidbeckingsale @rhornung67 can you take a look at this?

co-authored-by: Christian Trott <[email protected]> co-authored-by: David Zoeller Poliakoff <[email protected]> Adds optional Kokkos dependency and Kokkos as a submodule

co-authored-by: Christian Trott <[email protected]> co-authored-by: David Zoeller Poliakoff <[email protected]>

co-authored-by: Christian Trott <[email protected]> co-authored-by: David Zoeller Poliakoff <[email protected]> Note: This makes it so that one only compiles the basic kernel group when Kokkos is enabled. This required a few judicous ifdefs in some places in common. Eventually when the Kokkos variants for all the kernel groups are merged, those ifdefs are not needed.

…e format This commit effectively removes clang formatting inappropriately applied to these files.

…rmatted file This commit effectively reverts clang formatting applied to this file in later commits, thus minimizing changes to the original file.

Accept this change Co-authored-by: Jason Burmark <[email protected]>

Co-authored-by: Jason Burmark <[email protected]>

Suggestion accepted Co-authored-by: Jason Burmark <[email protected]>

src/common/RAJAPerfSuite.hpp

src/RAJAPerfSuiteDriver.cpp

src/CMakeLists.txt

artv3 · 2022-06-22T15:25:14Z

I'm having some trouble configuring with this PR.
I am using cmake version: 3.23.1

and I am getting the following error:

RAJAPerf/build/tpl/kokkos/compile_tests/CMakeFiles/CMakeTmp/CMakeLists.txt:
  Target "cmTC_5076b" requires the language dialect "CXX17" .  But the
  current compiler "GNU" does not support this, or CMake does not know the
  flags to enable it.
  
CMake Error at tpl/kokkos/cmake/kokkos_functions.cmake:917 (TRY_COMPILE):
  Failed to generate test project build system.
Call Stack (most recent call first):
  tpl/kokkos/cmake/kokkos_functions.cmake:937 (KOKKOS_CXX_COMPILER_CUDA_TEST)
  tpl/kokkos/cmake/kokkos_compiler_id.cmake:65 (kokkos_compilation)
  tpl/kokkos/cmake/kokkos_tribits.cmake:204 (INCLUDE)
  tpl/kokkos/CMakeLists.txt:170 (KOKKOS_SETUP_BUILD_ENVIRONMENT)

artv3 · 2022-06-22T15:41:10Z

I was able to get to configure by using a newer GNU - 8.3.1, but I then get warnings of the following type:

CMake Warning (dev) in test/CMakeLists.txt:
  Policy CMP0104 is not set: CMAKE_CUDA_ARCHITECTURES now detected for NVCC,
  empty CUDA_ARCHITECTURES not allowed.  Run "cmake --help-policy CMP0104"
  for policy details.  Use the cmake_policy command to set the policy and
  suppress this warning

then I get build errors:

RAJAPerf/tpl/kokkos/core/src/Cuda/Kokkos_Cuda_Half.hpp(667): error: more than one conversion function from "Kokkos::Experimental::half_t::impl_type" to a built-in type applies:
            function "__half::operator float() const"
            function "__half::operator short() const"
            function "__half::operator unsigned short() const"
            function "__half::operator int() const"
            function "__half::operator unsigned int() const"
            function "__half::operator long long() const"
            function "__half::operator unsigned long long() const"
            function "__half::operator __nv_bool() const"

ajpowelsnl · 2022-06-22T16:27:05Z

I'm having some trouble configuring with this PR. I am using cmake version: 3.23.1

and I am getting the following error:

RAJAPerf/build/tpl/kokkos/compile_tests/CMakeFiles/CMakeTmp/CMakeLists.txt:
  Target "cmTC_5076b" requires the language dialect "CXX17" .  But the
  current compiler "GNU" does not support this, or CMake does not know the
  flags to enable it.
  
CMake Error at tpl/kokkos/cmake/kokkos_functions.cmake:917 (TRY_COMPILE):
  Failed to generate test project build system.
Call Stack (most recent call first):
  tpl/kokkos/cmake/kokkos_functions.cmake:937 (KOKKOS_CXX_COMPILER_CUDA_TEST)
  tpl/kokkos/cmake/kokkos_compiler_id.cmake:65 (kokkos_compilation)
  tpl/kokkos/cmake/kokkos_tribits.cmake:204 (INCLUDE)
  tpl/kokkos/CMakeLists.txt:170 (KOKKOS_SETUP_BUILD_ENVIRONMENT)

Hmmm... Thanks for bringing this up, Arturo. Let me look into this. I was using gcc/9.1 + cuda/11.2, and it was OK for me. Have you updated your Kokkos submodule recently, i.e., at the project root , git submodule update --recursive --init ? This smells like an out-of-date Kokkos dependency (in RAJAPerf/tpl/kokkos). I will check builds with gcc/8.3.1.

artv3 · 2022-06-22T16:32:36Z

I'm having some trouble configuring with this PR. I am using cmake version: 3.23.1
and I am getting the following error:
RAJAPerf/build/tpl/kokkos/compile_tests/CMakeFiles/CMakeTmp/CMakeLists.txt:
  Target "cmTC_5076b" requires the language dialect "CXX17" .  But the
  current compiler "GNU" does not support this, or CMake does not know the
  flags to enable it.
  
CMake Error at tpl/kokkos/cmake/kokkos_functions.cmake:917 (TRY_COMPILE):
  Failed to generate test project build system.
Call Stack (most recent call first):
  tpl/kokkos/cmake/kokkos_functions.cmake:937 (KOKKOS_CXX_COMPILER_CUDA_TEST)
  tpl/kokkos/cmake/kokkos_compiler_id.cmake:65 (kokkos_compilation)
  tpl/kokkos/cmake/kokkos_tribits.cmake:204 (INCLUDE)
  tpl/kokkos/CMakeLists.txt:170 (KOKKOS_SETUP_BUILD_ENVIRONMENT)  
  
Hmmm... Thanks for bringing this up, Arturo. Let me look into this. I was using gcc/9.1 + cuda/11.2, and it was OK for me. Have you updated your Kokkos submodule recently, i.e., at the project root , git submodule update --recursive --init ? This smells like an out-of-date Kokkos dependency (in RAJAPerf/tpl/kokkos). I will check builds with gcc/8.3.1.

Yes I tried: git submodule update --recursive --init no luck on my end.

I am using cuda-11.0.2

ajpowelsnl · 2022-06-22T16:38:02Z

I'm having some trouble configuring with this PR. I am using cmake version: 3.23.1
and I am getting the following error:
RAJAPerf/build/tpl/kokkos/compile_tests/CMakeFiles/CMakeTmp/CMakeLists.txt:
  Target "cmTC_5076b" requires the language dialect "CXX17" .  But the
  current compiler "GNU" does not support this, or CMake does not know the
  flags to enable it.
  
CMake Error at tpl/kokkos/cmake/kokkos_functions.cmake:917 (TRY_COMPILE):
  Failed to generate test project build system.
Call Stack (most recent call first):
  tpl/kokkos/cmake/kokkos_functions.cmake:937 (KOKKOS_CXX_COMPILER_CUDA_TEST)
  tpl/kokkos/cmake/kokkos_compiler_id.cmake:65 (kokkos_compilation)
  tpl/kokkos/cmake/kokkos_tribits.cmake:204 (INCLUDE)
  tpl/kokkos/CMakeLists.txt:170 (KOKKOS_SETUP_BUILD_ENVIRONMENT)  
  
Hmmm... Thanks for bringing this up, Arturo. Let me look into this. I was using gcc/9.1 + cuda/11.2, and it was OK for me. Have you updated your Kokkos submodule recently, i.e., at the project root , git submodule update --recursive --init ? This smells like an out-of-date Kokkos dependency (in RAJAPerf/tpl/kokkos). I will check builds with gcc/8.3.1.
Yes I tried: git submodule update --recursive --init no luck on my end.

I am using cuda-11.0.2

@artv3 -- Here's the configuration I'm using on an NVIDIA server:

# Kokkos Cuda build
cmake \
-DENABLE_KOKKOS=ON \
-DRAJA_PERFSUITE_ENABLE_TESTS=OFF \
-DCUDA_ARCH=sm_70 \
-DENABLE_CUDA=ON \
-DKokkos_ARCH_VOLTA70=ON \
-DCMAKE_CUDA_ARCHITECTURES=70 \
-DCMAKE_BUILD_TYPE=Release .. \

Can you try this configuration? Again, this is with gcc/9.1 + CUDA/11.2

artv3 · 2022-06-22T16:57:59Z

Thanks, between that and switching the kokkos branch to master I was able to build. The hash for kokkos I got from git submodule was 2834f94af which gave me the following build error:

src/basic-kokkos/REDUCE3_INT-Kokkos.cpp(58): error: namespace "Kokkos::Experimental" has no member "min"

Running the suite via: ./bin/raja-perf.exe -k Basic_DAXPY
Gives me the following output for the timing-average.csv:

Mean Runtime Report (sec.)  ,  ,  ,  ,  ,  ,  ,                                                                                                                              
Kernel       , Base_Seq , Lambda_Seq , RAJA_Seq , Base_CUDA , Lambda_CUDA , RAJA_CUDA , Kokkos_Lambda                                                                        
Kernel       , default  , default    , default  , block_256 , block_256   , block_256 , default                                                                              
Basic_DAXPY  , 0.303750 ,   0.307818 , 0.309596 ,  0.016320 ,    0.016331 ,  0.016322 ,      0.016293

Based on run-time it looks like kokkos is running with the CUDA backend. For a fair comparison, it would be nice if the number of threads per block were reported.

artv3 · 2022-06-22T16:59:11Z

Could the following be added to a script?

# Kokkos Cuda build
cmake \
-DENABLE_KOKKOS=ON \
-DRAJA_PERFSUITE_ENABLE_TESTS=OFF \
-DCUDA_ARCH=sm_70 \
-DENABLE_CUDA=ON \
-DKokkos_ARCH_VOLTA70=ON \
-DCMAKE_CUDA_ARCHITECTURES=70 \
-DCMAKE_BUILD_TYPE=Release .. \

ajpowelsnl · 2022-06-22T17:51:13Z

Thanks, between that and switching the kokkos branch to master I was able to build. The hash for kokkos I got from git submodule was 2834f94af which gave me the following build error:

src/basic-kokkos/REDUCE3_INT-Kokkos.cpp(58): error: namespace "Kokkos::Experimental" has no member "min"

Running the suite via: ./bin/raja-perf.exe -k Basic_DAXPY Gives me the following output for the timing-average.csv:
Mean Runtime Report (sec.)  ,  ,  ,  ,  ,  ,  ,                                                                                                                              
Kernel       , Base_Seq , Lambda_Seq , RAJA_Seq , Base_CUDA , Lambda_CUDA , RAJA_CUDA , Kokkos_Lambda                                                                        
Kernel       , default  , default    , default  , block_256 , block_256   , block_256 , default                                                                              
Basic_DAXPY  , 0.303750 ,   0.307818 , 0.309596 ,  0.016320 ,    0.016331 ,  0.016322 ,      0.016293  
Based on run-time it looks like kokkos is running with the CUDA backend. For a fair comparison, it would be nice if the number of threads per block were reported.

Yes, I did most of the development on a CUDA machine. I will speak with @crtrott about the possibility of reporting the number of threads (per backend) block. Would it be possible to merge this PR now? We are going on a year now. If this work is not merged soon, then I will not be able to contribute the other Kokkos variants.

artv3 · 2022-06-22T17:53:44Z

Thanks, between that and switching the kokkos branch to master I was able to build. The hash for kokkos I got from git submodule was 2834f94af which gave me the following build error:
src/basic-kokkos/REDUCE3_INT-Kokkos.cpp(58): error: namespace "Kokkos::Experimental" has no member "min"
Running the suite via: ./bin/raja-perf.exe -k Basic_DAXPY Gives me the following output for the timing-average.csv:
Mean Runtime Report (sec.)  ,  ,  ,  ,  ,  ,  ,                                                                                                                              
Kernel       , Base_Seq , Lambda_Seq , RAJA_Seq , Base_CUDA , Lambda_CUDA , RAJA_CUDA , Kokkos_Lambda                                                                        
Kernel       , default  , default    , default  , block_256 , block_256   , block_256 , default                                                                              
Basic_DAXPY  , 0.303750 ,   0.307818 , 0.309596 ,  0.016320 ,    0.016331 ,  0.016322 ,      0.016293  
Based on run-time it looks like kokkos is running with the CUDA backend. For a fair comparison, it would be nice if the number of threads per block were reported.
Yes, I did most of the development on a CUDA machine. I will speak with @crtrott about the possibility of reporting the number of threads (per backend) block. Would it be possible to merge this PR now? We are going on a year now. If this work is not merged soon, then I will not be able to contribute the other Kokkos variants.

Approved! #248 (review)

ajpowelsnl · 2022-06-22T17:56:36Z

@artv3 I am looking into this error:

src/basic-kokkos/REDUCE3_INT-Kokkos.cpp(58): error: namespace "Kokkos::Experimental" has no member "min"

We recently moved min, max, etc. to the Kokkos namespace (and out of Kokkos::Experimental).

ajpowelsnl · 2022-06-22T17:57:11Z

Could the following be added to a script?

# Kokkos Cuda build
cmake \
-DENABLE_KOKKOS=ON \
-DRAJA_PERFSUITE_ENABLE_TESTS=OFF \
-DCUDA_ARCH=sm_70 \
-DENABLE_CUDA=ON \
-DKokkos_ARCH_VOLTA70=ON \
-DCMAKE_CUDA_ARCHITECTURES=70 \
-DCMAKE_BUILD_TYPE=Release .. \

Sure, no problem. I will put it in the scripts directory.

artv3 · 2022-06-22T18:04:53Z

@artv3 I am looking into this error:
src/basic-kokkos/REDUCE3_INT-Kokkos.cpp(58): error: namespace "Kokkos::Experimental" has no member "min" 
We recently moved min, max, etc. to the Kokkos namespace (and out of Kokkos::Experimental).

Yes, I think the kokkos submodule has to be updated. It currently points at hash: 2834f94af9b01debf67c1aaa3f0eb0c903d72c8d

artv3 · 2022-06-22T18:06:40Z

@artv3 I am looking into this error:
src/basic-kokkos/REDUCE3_INT-Kokkos.cpp(58): error: namespace "Kokkos::Experimental" has no member "min" 
We recently moved min, max, etc. to the Kokkos namespace (and out of Kokkos::Experimental).
Yes, I think the kokkos submodule has to be updated. It currently points at hash: 2834f94af9b01debf67c1aaa3f0eb0c903d72c8d

One way to do this is to go to the kokkos tpl folder and checkout the correct version of kokkos, commit and push the change. Using master worked for me, but perhaps we should use a released version.

davidbeckingsale requested review from davidbeckingsale and rhornung67 January 25, 2022 18:44