Fix the wrong type and pass real-number value with device_type to devices

[yhmtsai]

I face some type issue when I develop the half precision computation in ginkgo. I extract some fix here because the current way I tried in half precision computation may not be workable in the end.

There are two main fix in this PR:

1. the tuple type mismatch of indextype and valuetype
2. passing the remove_complex<ValueType> with as_device_type because we will use __half from vendor library not gko::half

[yhmtsai]

rebase!

[thoasm]

LGTM!
I would like to see a PR description change to detail why you now need to change the type between devices for remove_complex<T> floating point types.

[thoasm]

Do we actually need that? As far as I can see, gko::half should be the same on CPU vs. GPU.
Are you changing the type so that gko::half is a different type on CPU and GPU?
Looks like that is your plan in #1257.

[yhmtsai]

Yes, I will use __half on the GPU but gko::half on CPU.

[upsj]

LGTM!

[pratikvn]

There are a few extraneous as_device_type in the following kernels:

1. csr_kernels.cu: 302, 304, 315, 317 and maybe a few other places in this file
2. cb_gmres_kernels.cu: 351, 414

Probably the same in the hip kernels as well.

I think these places below are also missing a as_device_type ?:

1. idr_kernels.cu: 102

Does cublas, curand and cusparse have half precision support ? We dont seem to currently use as_device_type while passing the pointers to those wrappers.

Otherwise LGTM!

[yhmtsai]

curandGenerateNormal from cuda does not support half.
Some of cublas and cusparse might have half support. it will use as_culib_type

[thoasm]

LGTM!

[yhmtsai]

@pratikvn for idr curand, I do not convert it to device now because we only use the float double and generate 2x for complex (currently based on std::complex)

[sonarqubecloud]

Kudos, SonarCloud Quality Gate passed!   

0 Bugs
0 Vulnerabilities
0 Security Hotspots
0 Code Smells

No Coverage information
10.7% Duplication