Kokkos::deep_copy memory access failures

[mndevec]

This is caught for some of the SpGEMM tests described here: https://github.com/kokkos/kokkos-kernels/wiki/SpGEMM_Benchmarks. (For example coPaperDBLP).

Something for Kokkos::deep_copy between Kokkos versions 2.5 and 2.6 seems to introduce illegal memory accesses. This made some of the spgemm tests fail with some memory access issues.

I can replicate this issue in the below small test code on P100 with cuda 8.0 and gcc 5.3.

#include "Kokkos_Core.hpp"
#include <iostream>

int main() {

  Kokkos::initialize();
  std::cout << "Allocating view" << std::endl;
  Kokkos::View<int *, Kokkos::Cuda> very_big_view("test", 1960173568);
  Kokkos::fence();
  std::cout << "Deep Copy" << std::endl;

  Kokkos::deep_copy(very_big_view, 4);
  Kokkos::fence();
  std::cout << "Deep Copy Done" << std::endl;
  Kokkos::finalize();
}

The output is as below:

bash-4.2$ ./test.exe 
Allocating view
Deep Copy
terminate called after throwing an instance of 'std::runtime_error'
  what():  cudaDeviceSynchronize() error( cudaErrorIllegalAddress): an illegal memory access was encountered /ascldap/users/mndevec/shepard_work/kokkos/core/src/Cuda/Kokkos_Cuda_Impl.cpp:119
Traceback functionality not available

Aborted (core dumped)

@srajama1 @crtrott @ibaned @dsunder @ndellingwood

[crtrott]

Oh my ...
I am not sure how this slipped through our testing it really seems straight forward enough.

[srajama1]

I think we should add the performance testing framework developed by Michel so Kokkoskernels can stress test Kokkos develop once a week or so. This showed up because we were trying to reproduce an earlier run w/ large matrices. Such a test would have caught other performance problems too.

[crtrott]

I am also at somewhat of a loss what is going wrong. I tracked this down to this addition

  typedef   ViewFillA<Kokkos::View<int*, Kokkos::LayoutRight,
                                      Kokkos::Device<Kokkos::Cuda, Kokkos::CudaSpace>,
                                      Kokkos::MemoryTraits<1u> >,
                         Kokkos::LayoutLeft, Kokkos::Cuda, 1, int> FillFunctorType;
  FillFunctorType f(very_big_view,4);
  typedef Kokkos::RangePolicy<Kokkos::Cuda,Kokkos::IndexType<int>> policy_type;
  Kokkos::parallel_for("Copy2",policy_type(0,very_big_view.extent(0)), f);

This fails. But if I use int64_t as IndexType than it works.

[crtrott]

So this was somewhat bad luck. The code works if you use a length larger than INT_MAX and it works for everything smaller than INT_MAX/2 it fails for numbers in between. We have tests which exercise more than 2B but we don't have anything in the region 1B-2B which would have triggered this issue ...

[srajama1]

What is special about 1B-2B ?

[crtrott]

Maybe there is somewhere a *2 in the code path which overflows? I really don't know. I just tested whats going on to track down where this is going wrong and why we didn't catch it.

[srajama1]

Got it ! I was thinking of a * 2 as well but cannot see why that is needed for deep copy. Something weird.

[crtrott]

AH I think I know whats going on. Its based on how the internal loop structure in the CUDA impl works. Basically it does a loop with i + stride < end .
Stride in this case would actually be 1.xB and i could also be 1.xB so together they overflow.

[ndellingwood]

PR #1588 merged, marking this as InDevelop.