ICODE-ADC/lammps
4
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Update Kokkos library in LAMMPS to v3.2

Browse Source
This commit is contained in:
Stan Moore
2020-08-25 20:21:48 -06:00
parent 450fd12d31
commit 4d90c2b74b
1410 changed files with 19364 additions and 71953 deletions
Download Patch File Download Diff File Expand all files Collapse all files

147
lib/kokkos/core/unit_test/cuda/TestCuda_InterOp_Streams.cpp
View File

@ -42,103 +42,10 @@
//@HEADER
*/
#include <Kokkos_Core.hpp>
#include <cuda/TestCuda_Category.hpp>
#include <Test_InterOp_Streams.hpp>
namespace Test {
__global__ void offset_streams(int* p) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < 100) {
p[idx] += idx;
}
}
namespace {
struct FunctorRange {
Kokkos::View<int*, Kokkos::CudaSpace, Kokkos::MemoryTraits<Kokkos::Unmanaged>>
a;
FunctorRange(Kokkos::View<int*, Kokkos::CudaSpace,
Kokkos::MemoryTraits<Kokkos::Unmanaged>>
a_)
: a(a_) {}
KOKKOS_INLINE_FUNCTION
void operator()(const int i) const { a(i) += 1; }
};
struct FunctorRangeReduce {
Kokkos::View<int*, Kokkos::CudaSpace, Kokkos::MemoryTraits<Kokkos::Unmanaged>>
a;
FunctorRangeReduce(Kokkos::View<int*, Kokkos::CudaSpace,
Kokkos::MemoryTraits<Kokkos::Unmanaged>>
a_)
: a(a_) {}
KOKKOS_INLINE_FUNCTION
void operator()(const int i, int& lsum) const { lsum += a(i); }
};
struct FunctorMDRange {
Kokkos::View<int*, Kokkos::CudaSpace, Kokkos::MemoryTraits<Kokkos::Unmanaged>>
a;
FunctorMDRange(Kokkos::View<int*, Kokkos::CudaSpace,
Kokkos::MemoryTraits<Kokkos::Unmanaged>>
a_)
: a(a_) {}
KOKKOS_INLINE_FUNCTION
void operator()(const int i, const int j) const { a(i * 10 + j) += 1; }
};
struct FunctorMDRangeReduce {
Kokkos::View<int*, Kokkos::CudaSpace, Kokkos::MemoryTraits<Kokkos::Unmanaged>>
a;
FunctorMDRangeReduce(Kokkos::View<int*, Kokkos::CudaSpace,
Kokkos::MemoryTraits<Kokkos::Unmanaged>>
a_)
: a(a_) {}
KOKKOS_INLINE_FUNCTION
void operator()(const int i, const int j, int& lsum) const {
lsum += a(i * 10 + j);
}
};
struct FunctorTeam {
Kokkos::View<int*, Kokkos::CudaSpace, Kokkos::MemoryTraits<Kokkos::Unmanaged>>
a;
FunctorTeam(Kokkos::View<int*, Kokkos::CudaSpace,
Kokkos::MemoryTraits<Kokkos::Unmanaged>>
a_)
: a(a_) {}
KOKKOS_INLINE_FUNCTION
void operator()(
const Kokkos::TeamPolicy<Kokkos::Cuda>::member_type& team) const {
int i = team.league_rank();
Kokkos::parallel_for(Kokkos::TeamThreadRange(team, 10),
[&](const int j) { a(i * 10 + j) += 1; });
}
};
struct FunctorTeamReduce {
Kokkos::View<int*, Kokkos::CudaSpace, Kokkos::MemoryTraits<Kokkos::Unmanaged>>
a;
FunctorTeamReduce(Kokkos::View<int*, Kokkos::CudaSpace,
Kokkos::MemoryTraits<Kokkos::Unmanaged>>
a_)
: a(a_) {}
KOKKOS_INLINE_FUNCTION
void operator()(const Kokkos::TeamPolicy<Kokkos::Cuda>::member_type& team,
int& lsum) const {
int i = team.league_rank();
int team_sum;
Kokkos::parallel_reduce(
Kokkos::TeamThreadRange(team, 10),
[&](const int j, int& tsum) { tsum += a(i * 10 + j); }, team_sum);
Kokkos::single(Kokkos::PerTeam(team), [&]() { lsum += team_sum; });
}
};
} // namespace
// Test Interoperability with Cuda Streams
TEST(cuda, raw_cuda_streams) {
cudaStream_t stream;
@ -147,45 +54,47 @@ TEST(cuda, raw_cuda_streams) {
Kokkos::initialize(arguments);
int* p;
cudaMalloc(&p, sizeof(int) * 100);
using MemorySpace = typename TEST_EXECSPACE::memory_space;
{
Kokkos::Cuda cuda0(stream);
Kokkos::View<int*, Kokkos::CudaSpace> v(p, 100);
Kokkos::deep_copy(cuda0, v, 5);
TEST_EXECSPACE space0(stream);
Kokkos::View<int*, TEST_EXECSPACE> v(p, 100);
Kokkos::deep_copy(space0, v, 5);
int sum;
Kokkos::parallel_for("Test::cuda::raw_cuda_stream::Range",
Kokkos::RangePolicy<Kokkos::Cuda>(cuda0, 0, 100),
FunctorRange(v));
Kokkos::RangePolicy<TEST_EXECSPACE>(space0, 0, 100),
FunctorRange<MemorySpace>(v));
Kokkos::parallel_reduce(
"Test::cuda::raw_cuda_stream::RangeReduce",
Kokkos::RangePolicy<Kokkos::Cuda, Kokkos::LaunchBounds<128, 2>>(cuda0,
0, 100),
FunctorRangeReduce(v), sum);
cuda0.fence();
Kokkos::RangePolicy<TEST_EXECSPACE, Kokkos::LaunchBounds<128, 2>>(
space0, 0, 100),
FunctorRangeReduce<MemorySpace>(v), sum);
space0.fence();
ASSERT_EQ(600, sum);
Kokkos::parallel_for("Test::cuda::raw_cuda_stream::MDRange",
Kokkos::MDRangePolicy<Kokkos::Cuda, Kokkos::Rank<2>>(
cuda0, {0, 0}, {10, 10}),
FunctorMDRange(v));
Kokkos::parallel_reduce("Test::cuda::raw_cuda_stream::MDRangeReduce",
Kokkos::MDRangePolicy<Kokkos::Cuda, Kokkos::Rank<2>,
Kokkos::LaunchBounds<128, 2>>(
cuda0, {0, 0}, {10, 10}),
FunctorMDRangeReduce(v), sum);
cuda0.fence();
Kokkos::MDRangePolicy<TEST_EXECSPACE, Kokkos::Rank<2>>(
space0, {0, 0}, {10, 10}),
FunctorMDRange<MemorySpace>(v));
Kokkos::parallel_reduce(
"Test::cuda::raw_cuda_stream::MDRangeReduce",
Kokkos::MDRangePolicy<TEST_EXECSPACE, Kokkos::Rank<2>,
Kokkos::LaunchBounds<128, 2>>(space0, {0, 0},
{10, 10}),
FunctorMDRangeReduce<MemorySpace>(v), sum);
space0.fence();
ASSERT_EQ(700, sum);
Kokkos::parallel_for("Test::cuda::raw_cuda_stream::Team",
Kokkos::TeamPolicy<Kokkos::Cuda>(cuda0, 10, 10),
FunctorTeam(v));
Kokkos::TeamPolicy<TEST_EXECSPACE>(space0, 10, 10),
FunctorTeam<MemorySpace, TEST_EXECSPACE>(v));
Kokkos::parallel_reduce(
"Test::cuda::raw_cuda_stream::Team",
Kokkos::TeamPolicy<Kokkos::Cuda, Kokkos::LaunchBounds<128, 2>>(cuda0,
10, 10),
FunctorTeamReduce(v), sum);
cuda0.fence();
Kokkos::TeamPolicy<TEST_EXECSPACE, Kokkos::LaunchBounds<128, 2>>(
space0, 10, 10),
FunctorTeamReduce<MemorySpace, TEST_EXECSPACE>(v), sum);
space0.fence();
ASSERT_EQ(800, sum);
}
Kokkos::finalize();
@ -193,7 +102,7 @@ TEST(cuda, raw_cuda_streams) {
CUDA_SAFE_CALL(cudaDeviceSynchronize());
cudaStreamDestroy(stream);
int* h_p = new int[100];
int h_p[100];
cudaMemcpy(h_p, p, sizeof(int) * 100, cudaMemcpyDefault);
CUDA_SAFE_CALL(cudaDeviceSynchronize());
int64_t sum = 0;

1
lib/kokkos/core/unit_test/cuda/TestCuda_Other.cpp
View File

@ -48,7 +48,6 @@
#include <TestAggregate.hpp>
#include <TestMemoryPool.hpp>
#include <TestCXX11.hpp>
#include <TestTile.hpp>
#include <TestViewCtorPropEmbeddedDim.hpp>
#include <TestViewLayoutTiled.hpp>

1
lib/kokkos/core/unit_test/cuda/TestCuda_ViewLayoutStrideAssignment.cpp
View File

@ -44,3 +44,4 @@
#include <cuda/TestCuda_Category.hpp>
#include <TestViewLayoutStrideAssignment.hpp>
#include <TestIrregularLayout.hpp>