rocSPARSE 3.0.2 for ROCm 6.0 breaking changes

[maartenarnst]

It appears that rocSPARSE 3.0.2 that comes with ROCm 6.0 makes changes that break things in Kokkos-kernels.

As compared with ROCm 5.7, the file rocm_version.h is now moved to a subfolder rocm-core of ROCm's include dir. Potential solutions could be either including this folder as an include directory or modifying lines

• kokkos-kernels/sparse/src/KokkosSparse_Utils_rocsparse.hpp

  Line 24 in f429f6e

  #include <rocm_version.h>

• kokkos-kernels/sparse/tpls/KokkosSparse_spmv_tpl_spec_decl.hpp

  Line 361 in f429f6e

  #include <rocm_version.h>

A more challenging change is that certain enumerators for algorithms have been changed. For instance, rocsparse_spmv_stage_auto is no longer defined. And so it seems it becomes necessary to specify explicitly the stages in lines

• kokkos-kernels/sparse/tpls/KokkosSparse_spmv_tpl_spec_decl.hpp

  Lines 444 to 451 in f429f6e

  	#if KOKKOSSPARSE_IMPL_ROCM_VERSION >= 50400
  	KOKKOS_ROCSPARSE_SAFE_CALL_IMPL(rocsparse_spmv_ex(
  	handle, myRocsparseOperation, &alpha, Aspmat, vecX, &beta, vecY,
  	compute_type, alg, rocsparse_spmv_stage_auto, &buffer_size, tmp_buffer));
  	KOKKOS_IMPL_HIP_SAFE_CALL(hipMalloc(&tmp_buffer, buffer_size));
  	KOKKOS_ROCSPARSE_SAFE_CALL_IMPL(rocsparse_spmv_ex(
  	handle, myRocsparseOperation, &alpha, Aspmat, vecX, &beta, vecY,
  	compute_type, alg, rocsparse_spmv_stage_auto, &buffer_size, tmp_buffer));

but these stages appear to depend on the algorithm.

Compilation ends with an error on these lines involving rocsparse_spmv_stage_auto . I am unsure how to solve the issue, and so I haven't been able to try out the compilation beyond this problem. It's possible there might be more issues to solve to make things work with the updated ROCm.

[maartenarnst]

It seems these issues have already been solved in the current develop branch. (I had observed these issues when compiling kokkos-kernels as part of Trilinos).

It may be worth noting that the example on

-https://rocm.docs.amd.com/projects/rocSPARSE/en/latest/generic.html#_CPPv417rocsparse_spmv_ex16rocsparse_handle19rocsparse_operationPKvK21rocsparse_spmat_descrK21rocsparse_dnvec_descrPKvK21rocsparse_dnvec_descr18rocsparse_datatype18rocsparse_spmv_alg20rocsparse_spmv_stageP6size_tPv

goes through all three stages rocsparse_spsv_stage_buffer_size, rocsparse_spsv_stage_preprocess, rocsparse_spsv_stage_compute.

The current fix

• kokkos-kernels/sparse/tpls/KokkosSparse_spmv_tpl_spec_decl.hpp

  Lines 444 to 453 in 4315f9a

  	#if KOKKOSSPARSE_IMPL_ROCM_VERSION >= 60000
  	KOKKOS_ROCSPARSE_SAFE_CALL_IMPL(
  	rocsparse_spmv(handle, myRocsparseOperation, &alpha, Aspmat, vecX, &beta,
  	vecY, compute_type, alg, rocsparse_spmv_stage_buffer_size,
  	&buffer_size, tmp_buffer));
  	KOKKOS_IMPL_HIP_SAFE_CALL(hipMalloc(&tmp_buffer, buffer_size));
  	KOKKOS_ROCSPARSE_SAFE_CALL_IMPL(
  	rocsparse_spmv(handle, myRocsparseOperation, &alpha, Aspmat, vecX, &beta,
  	vecY, compute_type, alg, rocsparse_spmv_stage_compute,
  	&buffer_size, tmp_buffer));

omits the rocsparse_spmv_stage_preprocess stage. While this may be appropriate for the "stream" algorithm, it's not clear that omitting the preprocessing step is appropriate for the "default" algorithm.

• kokkos-kernels/sparse/tpls/KokkosSparse_spmv_tpl_spec_decl.hpp

  Lines 436 to 442 in 4315f9a

  	if (controls.isParameter("algorithm")) {
  	const std::string algName = controls.getParameter("algorithm");
  	if (algName == "default")
  	alg = rocsparse_spmv_alg_default;
  	else if (algName == "merge")
  	alg = rocsparse_spmv_alg_csr_stream;
  	}

Also not that I've just introduced PR #2100 to fix a minor issue with compiling when both rocsparse and mkl are enabled.

@seanofthemillers @lucbv

[seanofthemillers]

On the Trilinos, these fixes are being incorporated in trilinos/Trilinos#12681 (holdover until a new KK snapshot goes in), though there are a couple other ROCm fixes going in that PR at the same time.

[ndellingwood]

The kokkos-kernels snapshot to Trilinos PR trilinos/Trilinos#12707 has merged as well as trilinos/Trilinos#12681 , thanks for the fixes @seanofthemillers !