Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

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

Merged
merged 2 commits into from
Feb 9, 2023
Merged

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

merged 2 commits into from
Feb 9, 2023

Conversation

yhmtsai
Copy link
Member

@yhmtsai yhmtsai commented Jan 5, 2023 •
edited
Loading

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 yhmtsai self-assigned this Jan 5, 2023
@ginkgo-bot ginkgo-bot added mod:cuda This is related to the CUDA module. mod:hip This is related to the HIP module. type:factorization This is related to the Factorizations type:matrix-format This is related to the Matrix formats type:multigrid This is related to multigrid labels Jan 5, 2023
@yhmtsai
Copy link
Member Author

yhmtsai commented Jan 7, 2023

rebase!

@yhmtsai yhmtsai mentioned this pull request Jan 12, 2023
Open
12 tasks
thoasm
thoasm approved these changes Jan 13, 2023
View reviewed changes
Copy link
Member

@thoasm thoasm left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

cuda/factorization/par_ilut_filter_kernel.cu
old_row_ptrs, as_cuda_type(old_vals), num_rows, threshold,
new_row_ptrs, lower);
old_row_ptrs, as_cuda_type(old_vals), num_rows,
as_cuda_type(threshold), new_row_ptrs, lower);
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Copy link
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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

@yhmtsai yhmtsai added the 1:ST:ready-for-review This PR is ready for review label Feb 8, 2023
@yhmtsai yhmtsai requested review from a team and thoasm February 8, 2023 10:34
upsj
upsj approved these changes Feb 8, 2023
View reviewed changes
Copy link
Member

@upsj upsj left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM!

common/cuda_hip/matrix/fbcsr_kernels.hpp.inc Show resolved Hide resolved
pratikvn
pratikvn approved these changes Feb 8, 2023
View reviewed changes
Copy link
Member

@pratikvn pratikvn left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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!

common/unified/multigrid/pgm_kernels.cpp Outdated Show resolved Hide resolved
common/unified/multigrid/pgm_kernels.cpp Outdated Show resolved Hide resolved
@yhmtsai
Copy link
Member Author

yhmtsai commented Feb 8, 2023

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

thoasm
thoasm approved these changes Feb 8, 2023
View reviewed changes
Copy link
Member

@thoasm thoasm left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM!

@yhmtsai yhmtsai added 1:ST:ready-to-merge This PR is ready to merge. and removed 1:ST:ready-for-review This PR is ready for review labels Feb 9, 2023
@yhmtsai
Copy link
Member Author

yhmtsai commented Feb 9, 2023

@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)

@yhmtsai yhmtsai changed the title fix some type issue Fix the wrong type and pass real-number value with device_type to devices Feb 9, 2023
@yhmtsai @pratikvn
delete some redundent as_device_type
d7b41c7 
Co-authored-by: Pratik Nayak <pratikvn@protonmail.com>
@yhmtsai yhmtsai merged commit 19b7249 into develop Feb 9, 2023
@yhmtsai yhmtsai deleted the fixtype branch February 9, 2023 22:46
@sonarcloud
Copy link

sonarcloud bot commented Feb 10, 2023

Kudos, SonarCloud Quality Gate passed!    Quality Gate passed

Bug A 0 Bugs
Vulnerability A 0 Vulnerabilities
Security Hotspot A 0 Security Hotspots
Code Smell A 0 Code Smells

No Coverage information No Coverage information
10.7% 10.7% Duplication

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers

@upsj upsj upsj approved these changes

@pratikvn pratikvn pratikvn approved these changes

@thoasm thoasm thoasm approved these changes

Assignees

@yhmtsai yhmtsai

Labels
1:ST:ready-to-merge This PR is ready to merge. mod:cuda This is related to the CUDA module. mod:hip This is related to the HIP module. type:factorization This is related to the Factorizations type:matrix-format This is related to the Matrix formats type:multigrid This is related to multigrid
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.
5 participants
@yhmtsai @upsj @pratikvn @thoasm @ginkgo-bot