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.

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

Fixes dpctl.tensor.round on CUDA devices #1700

Merged
merged 1 commit into from
May 29, 2024
Merged

Conversation

ndgrigorian
Copy link
Collaborator

When compiled for CUDA, std::rint would incorrectly round values halfway between two integers towards 0 (i.e., 1.5 -> 1.). The array API specification requires that these values be rounded to the nearest even integer instead.

To resolve this, std::rint has been replaced with sycl::rint, which does not rely on the current floating-point rounding mode (see SYCL specification).

As was pointed out at the time of implementation the floating-point rounding mode can vary between devices.

  • Have you provided a meaningful PR description?
  • Have you added a test, reproducer or referred to an issue with a reproducer?
  • Have you tested your changes locally for CPU and GPU devices?
  • Have you made sure that new changes do not introduce compiler warnings?
  • Have you checked performance impact of proposed changes?
  • If this PR is a work in progress, are you opening the PR as a draft?

…vidia hardware

When compiled for CUDA, `std::rint` would incorrectly round values halfway between two integers toward 0, rather than to the nearest even number as required per array API.

`sycl::rint` avoids such issues by not relying on the current rounding mode
Copy link

github-actions bot commented May 29, 2024

Deleted rendered PR docs from intelpython.github.com/dpctl, latest should be updated shortly. 🤞

Copy link

Array API standard conformance tests for dpctl=0.18.0dev0=py310h15de555_29 ran successfully.
Passed: 890
Failed: 11
Skipped: 91

Copy link
Collaborator

@oleksandr-pavlyk oleksandr-pavlyk left a comment

Choose a reason for hiding this comment

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

It might be a good idea to scan the code base for remaining uses of std namespace transcendental functions and replace those one by one too

@coveralls
Copy link
Collaborator

Coverage Status

coverage: 87.911%. remained the same
when pulling d8705a2 on fix-round-for-nvidia
into 1de00cb on master.

@ndgrigorian
Copy link
Collaborator Author

It might be a good idea to scan the code base for remaining uses of std namespace transcendental functions and replace those one by one too

I agree. I'll do this as a separate PR, but I think it's a good idea too.

@oleksandr-pavlyk
Copy link
Collaborator

I have opened gh-1701 for build break with LLVM SYCL compiler, it is unrelated to changes in this PR.

@ndgrigorian ndgrigorian merged commit 7ae4303 into master May 29, 2024
59 of 60 checks passed
@ndgrigorian ndgrigorian deleted the fix-round-for-nvidia branch May 29, 2024 22:24
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

3 participants