Page MenuHomePhabricator

[test-suite,CUDA] Add a test case to test the edge cases for the implementation of llvm.round intrinsic in the PTX backend.
ClosedPublic

Authored by bixia on Mar 28 2019, 11:07 AM.

Diff Detail

Repository
rL LLVM

Event Timeline

bixia created this revision.Mar 28 2019, 11:07 AM
Herald added a project: Restricted Project. · View Herald TranscriptMar 28 2019, 11:07 AM
bixia added a reviewer: tra.Mar 28 2019, 11:09 AM
tra added inline comments.Mar 28 2019, 3:21 PM
External/CUDA/round.cu
16 ↗(On Diff #192685)

I'd use a name different from the standard libm function. do_round perhaps?

17 ↗(On Diff #192685)

__builtin_roundf() takes float as an argument, so if T is double, it's not doing what you wanted it to do.
For double you need __builtin_round()

24–47 ↗(On Diff #192685)

This is way too convoluted. We do have assert available on device side.
Structuring the tests the way External/CUDA/cmath.cu does would be much simpler. E.g.:

__device__ void test_a(int x) {
   assert(__builtin_roundf(0.5f + x) == 1.0f);
   assert(__builtin_roundf(-0.5f + x) == -1.0f);
...
}

__global__ void tests(int x)
{
  test_a(x);
  ...
}
int main() {
   tests<<<1,1>>>(0);
   ...
}

Note: x is a coarse way to block constant evaluation. It may not be needed if the test is run unoptimized.

52–53 ↗(On Diff #192685)

It would be good to add some comments why particular values were chosen. My guess is that be the smallest value with lsb of mantissa representing less than 1.0 and thus should not be affected by rounding.

bixia updated this revision to Diff 192755.Mar 28 2019, 6:17 PM

Rename round.cu to test_round.cu.
Fix the intrinsic for double.
Change the test to use assert and a variable to prevent constant folding.

bixia updated this revision to Diff 192762.Mar 28 2019, 6:45 PM

Fix the test to pass in a scalar value instead of a pointer.
Remove unnecessary device memory allocation and copy.
Add a call to cudaDeviceSynchronize.

tra accepted this revision.Mar 28 2019, 7:33 PM
tra added inline comments.
External/CUDA/test_round.cu
37 ↗(On Diff #192762)

You may want to add a check for a similarly large negative value, too.

This revision is now accepted and ready to land.Mar 28 2019, 7:35 PM
bixia updated this revision to Diff 192876.Mar 29 2019, 11:38 AM

Modify test to also check a negative value beyond -max(float).

This revision was automatically updated to reflect the committed changes.