TPL Support for BLAS functions (nrm2, axpy, dot, gemm) using CuBLAS (Issue #247)

[vqd8a]

Added cuBLAS TPL Support for nrm2, axpy, dot. More to come.

[mhoemmen]

See comments. It's likely that this change will cause build warnings, due to signed / unsigned comparisons. Are you building with -Wall?

[mhoemmen]

Does this even work? wildcard should expand * etc., but there's no * here.

[vqd8a]

@mhoemmen I did not build with -Wall. But it is still fine with -Wall. wildcard works. I use it here just for other possible PRs. To avoid confusion, I will wildcard.

[vqd8a]

remove wildcard

[mhoemmen]

X.extent(0) may return size_t, which is usually unsigned. Comparison with INT_MAX may cause a build warning ("signed / unsigned comparison").

[mhoemmen]

This may cause a build warning, since it assigns size_t (usually unsigned) to int. Try this instead:

const int N = static_cast<int> (X.extent (0));

[mhoemmen]

See above.

[mhoemmen]

See above.

[mhoemmen]

See above.

[mhoemmen]

See above.

[mhoemmen]

See above.

[mhoemmen]

See above.

[mhoemmen]

See above.

[vqd8a]

Thanks @mhoemmen for the comments. I will fix these.

[srajama1]

Is the BLAS3 (GEMM especially) done ? We have an use for it, so it would be nice to prioritize that next if you can.

[srajama1]

Do we have to do anything in generate_makefile ? How about the integration w/ Trilinos TPL mechanisms ?

[vqd8a]

@srajama1 I made a minor change in Makefile.kokkos-kernels which adds the source file for handling cuBLAS.
I have not had an experience in Trilinos integration. Maybe @crtrott can answer.
I haven't done BLAS3 yet. But can work on it now.

[kyungjoo-kim]

I remembered that @mhoemmen commented that two versions of cublas should be harnessed for sierra build. @mhoemmen How do we test it ?

[mhoemmen]

@kyungjoo-kim Just don't expose any of the cublas headers in any header file, and you should be fine. The main issue is that FETI uses v2 of the cuBLAS API. If you include the v1 cuBLAS header, FETI won't build.

[mhoemmen]

@kyungjoo-kim It's OK to use both versions of the cuBLAS API in the same executable, but you aren't allowed to include both header files in the same compilation unit.

[mhoemmen]

@kyungjoo-kim If you must expose the header file (ask yourself why you really must; you almost certainly do not), then please use v2 of the cuBLAS API. That shouldn't break FETI.

[crtrott]

At minimum move a file

[crtrott]

This file should not be in a generated_specializations_cpp. That should only have auto generated files.

[vqd8a]

@crtrott I am moving this file to tpls

[crtrott]

Do we need to do this via a pointer or a create function or something. cublasCreate needs to be called after Kokkos::initialize().

[vqd8a]

@crtrott cublasCreate is already called after Kokkos::initialize()

[mhoemmen]

The singleton won't get destroyed until after main(). I'm not comfortable with the idea of CUDA state persisting after Kokkos::finalize. Why not get rid of the destructor, and have the constructor do the following?

CudaBlasSingleton::CudaBlasSingleton()
{    
  cublasCreate( & handle );
  cublasStatus_t stat = cublasCreate(&handle);
  if (stat != CUBLAS_STATUS_SUCCESS) {
    Kokkos::abort("CUBLAS initialization failed!");
  }
  Kokkos::push_finalize_hook ([&] () { cublasDestroy (handle); });
}

[mhoemmen]

Also, please never print anything (by default) unless something went wrong. Those printf statements will be very annoying if running on many nodes.

[crtrott]

v1 is deprecated. We should not include that ...

[vqd8a]

@srajama1 cuBLAS GEMM support has been added.

[vqd8a]

Changed PR name for partially addressing #247 for convenience. I will create separate PRs for the remaining BLAS functions.

[srajama1]

Is the handle that we expose for sparse cases an option here so users create the handle and destroy it both for sparse and dense KK functions ? Discussed this w/ @vqd8a already. @crtrott what do you think ? Can Tpetra create and delete the handles ?

[mhoemmen]

If the handle only ever gets created with the default CUDA stream, then the right way to deal with it is this:

1. Create it once, lazily, on demand
2. Add it as a Kokkos::finalize hook

If it's possible to create the handle with other CUDA streams, then users (such as Tpetra and other downstream software) need a way to create, manage, and destroy handles.

[srajama1]

I was wondering if the users like Tpetra would like to control things like cublasSetMathMode, cublasSetAtomicMode, cublasSetMatrix/VectorAsync, and cuBLASSetStream. KK can leave most of it to default, but making sure we don't force it. The disadvantage of the handle is of course the users have to worry about it even w/ other libraries like MKL (or with no libraries).

[mhoemmen]

Shouldn't you use appropriate kokkos-kernels macros to detect whether those Scalar types (e.g., float) are enabled? Otherwise, you'll be instantiating for types for which the user did not want to instantiate. This will increase build time and library size.

[mhoemmen]

The singleton won't get destroyed until after main(). I'm not comfortable with the idea of CUDA state persisting after Kokkos::finalize. Why not get rid of the destructor, and have the constructor do the following?

CudaBlasSingleton::CudaBlasSingleton()
{    
  cublasCreate( & handle );
  cublasStatus_t stat = cublasCreate(&handle);
  if (stat != CUBLAS_STATUS_SUCCESS) {
    Kokkos::abort("CUBLAS initialization failed!");
  }
  Kokkos::push_finalize_hook ([&] () { cublasDestroy (handle); });
}

[mhoemmen]

Also, please never print anything (by default) unless something went wrong. Those printf statements will be very annoying if running on many nodes.

[mhoemmen]

I was wondering if the users like Tpetra would like to control things like cublasSetMathMode, cublasSetAtomicMode, cublasSetMatrix/VectorAsync, and cuBLASSetStream.

• cublasSet{Matrix,Vector}Async appear to be redundant with respect to Kokkos::deep_copy
• "math mode" only appears relevant for float and half, not for double etc.
• "atomic mode" would give nondeterministic results so I'd rather not enable it unless somebody really insists

That leaves setting the stream. I could see users wanting to do that. MAGMA has a handle tied to a CUDA stream, for example. Better exposure of distinct CUDA streams would help us overlap data movement and computation.

The main value of hiding the stream is for gradually porting from existing interfaces like Teuchos::BLAS to this new interface. I think that's a good idea. It will be a while before Trilinos is ready to use multiple CUDA streams, so maybe we don't have to worry about that for now. Even three-argument Kokkos::deep_copy currently ignores its execution space instance argument: https://github.com/kokkos/kokkos/blob/d3a941925cbfb71785d8ea68259123ed52d3f9da/core/src/Kokkos_CopyViews.hpp#L1511

[mhoemmen]

@vqd8a Thanks for fixing the CudaBlasSingleton constructor! I appreciate also that you took out the printf statements :-) .

I'm dismissing my review because I'm not sure how kokkos-kernels wants to do ETI. Also, since this is a BLAS interface and we want to use it to replace things like Teuchos::BLAS, it may make sense to instantiate for the four BLAS types, independently of kokkos-kernels' ETI settings.

[vqd8a]

Thanks @mhoemmen for your suggestions. While waiting for decision, I have just fixed the use of kokkos-kernes macros to detect Scalar types as you suggested.

[mhoemmen]

Just a comment -- thanks!

[mhoemmen]

How come the complex double and complex float macros end with underscores, but the double and float ones don't?

[vqd8a]

I think we can use either without or with underscore, since in the KokkosKernels_config.h, it uses

#if defined(KOKKOSKERNELS_INST_COMPLEX_DOUBLE)
#define KOKKOSKERNELS_INST_KOKKOS_COMPLEX_DOUBLE_
#endif

but I just followed the convention in hpp files in \src\impl\generated_specializations_hpp

[mhoemmen]

OK, just checking :)

[srajama1]

@crtrott : Can you check this ? Need your opinion on two things, exposing (or not exposing) the handle and the ETI use case mentioned above.

[vqd8a]

@srajama1 , I talked to @crtrott last week. @crtrott said not exposing handles could be the option for now. and leave exposing for future. If it is the case, I can run the spot-check when the testbeds are back on.

[srajama1]

Agreed. Run the spot check when you can. @crtrott still needs to approve it as it is blocked on his request for changes.

[crtrott]

Looks good to me if you can run the spot-check. But also open a follow up: we may want to have the TPL variants independent of the INST macros. You see the INST macros don't say whether something is available, they say wether we do ETI. As long as you don't also define ETI only, all the other calls can still be made. And I think it is fine in that case to have the TPL calls available.

[vqd8a]

Spot check on white:
kokkos-kernels/scripts/test_all_sandia cuda --spot-check --with-cuda-options=enable_lambda --with-tpls=cublas

Going to test compilers:  cuda/8.0.44 cuda/9.0.103
Testing compiler cuda/8.0.44
  Starting job cuda-8.0.44-Cuda_OpenMP-release
  PASSED cuda-8.0.44-Cuda_OpenMP-release
  Starting job cuda-8.0.44-Cuda_Serial-release
  PASSED cuda-8.0.44-Cuda_Serial-release
Testing compiler cuda/9.0.103
  Starting job cuda-9.0.103-Cuda_OpenMP-release
  PASSED cuda-9.0.103-Cuda_OpenMP-release
  Starting job cuda-9.0.103-Cuda_Serial-release
  PASSED cuda-9.0.103-Cuda_Serial-release
#######################################################
PASSED TESTS
#######################################################
cuda-8.0.44-Cuda_OpenMP-release build_time=823 run_time=751
cuda-8.0.44-Cuda_Serial-release build_time=792 run_time=1078
cuda-9.0.103-Cuda_OpenMP-release build_time=862 run_time=761
cuda-9.0.103-Cuda_Serial-release build_time=800 run_time=1087

[srajama1]

Once you have a KNL spot check we can merge this.

[srajama1]

Ignore the silly request for KNL, this shouldn't affect any KNL, so merging.