Fix device-side conversions to 64-bit integer types. #1205
Conversation
|
!build |
|
CI MESSAGE: [873349]: BUILD STARTED |
|
CI MESSAGE: [873349]: BUILD FAILED |
|
!build |
|
CI MESSAGE: [873548]: BUILD STARTED |
|
!build |
|
CI MESSAGE: [873580]: BUILD STARTED |
|
CI MESSAGE: [873548]: BUILD PASSED |
|
CI MESSAGE: [873580]: BUILD PASSED |
| return __double2int_rn(f); | ||
| } | ||
| __device__ unsigned cuda_round_helper(double f, unsigned) { // NOLINT | ||
| inline __device__ unsigned cuda_round_helper(double f, unsigned) { // NOLINT | ||
| return __double2uint_rn(f); | ||
| } |
There was a problem hiding this comment.
would be good to add a comment explaining the reason behind the workaround, for future readers of this file
| @@ -139,6 +141,9 @@ inline __device__ DeviceString dev_to_string(long long x) { // NOLINT | |||
| if (x < 0) { | |||
| neg = true; | |||
| x = -x; | |||
| if (x < 0) { | |||
| return "-9223372036854775808"; | |||
There was a problem hiding this comment.
don't really understand this
There was a problem hiding this comment.
That's decimal representation of -2^63. U2 numbers are asymmetric and n-bit signed number cannot represent +2^(n-1), but can represent -2^(n-1). Negating -(-2^(n-1)) == -2^n(-1) due to overflow.
I could add a special case for this particular number.
Work around a bug in __float2ll_rn and related functions. Add tests for device-side conversions to 64-bit integers. Fix `dev_to_string(long long)` for input value of -2^63. Signed-off-by: Michal Zientkiewicz <michalz@nvidia.com>
|
!build |
|
CI MESSAGE: [875079]: BUILD STARTED |
|
CI MESSAGE: [875079]: BUILD PASSED |
Signed-off-by: Michal Zientkiewicz michalz@nvidia.com
Why we need this PR?
What happened in this PR?
JIRA TASK: [DALI-1031]