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

[tra]

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]

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]

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]

External/CUDA/test_round.cu

37

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

[bixia]

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

External/CUDA/CMakeLists.txt

macro(create_local_cuda_tests VariantSuffix)
   create_one_local_test(future future.cu)
   # We only need SIMD tests on CUDA-8.0 to verivy that our reference is correct
   # and matches NVIDIA-provided one.  and on CUDA-9.2 to verify that clang's
   # implementation matches the reference.  This test also happens to be the
   # longest one, so by not running unnecessary instances we speed up cuda
   # buildbot a lot.
   create_one_local_test_f(simd simd.cu
                           "cuda-(8[.]0|9[.]2)-c[+][+]11-libc[+][+]")
+  create_one_local_test(test_round test_round.cu)
 endmacro()
 
 macro(thrust_make_test_name TestName TestSourcePath)
   string(REPLACE "${THRUST_PATH}/testing/" "" _tmp ${TestSourcePath})
   string(REPLACE "/" "-" _tmp ${_tmp})
   string(REPLACE "<" "_" _tmp ${_tmp})
   string(REPLACE ">" "_" _tmp ${_tmp})
   string(REGEX REPLACE "\.(cpp|cu)$" "" _tmp ${_tmp})

External/CUDA/test_round.cu

+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <iostream>
+#include <cassert>
+
+// Test the implementation of llvm intrinsic round. In particular, when the
+// source is equidistant between two integers, it rounds away from zero.
+//
+// In CUDA libdevice, the implementation of round separates the values into
+// three regions and uses a region specific rounding method to calculate
+// the result:
+//    abs(x) <= 0.5
+//    2 ^ 23 > abs(x) > 0.5 (float)
+//    abs(x) >= 2 ^ 23 (float)
+//    For double, 2 ^ 23 above is replaced with 2 ^ 52
+//
+// The PTX backend implements round in a similar way. We chose the test values
+// based on this.
+
+__global__ void test_round(float v) {
+  assert(__builtin_roundf(-0.5f + v) == -1.0f);
+  assert(__builtin_roundf(8.5f + v) == 9.0f);
+  assert(__builtin_roundf(-8.38861e+06f + v) == -8.38861e+06f);
+  assert(__builtin_roundf(8.38861e+06f + v) == 8.38861e+06f);
+
+  assert(__builtin_round(0.5 + v) == 1.0f);
+  assert(__builtin_round(-8.5 + v) == -9.0f);
+  assert(__builtin_round(4.5035996e+15 + v) == 4.5035996e+15);
+  assert(__builtin_round(-4.5035996e+15 + v) == -4.5035996e+15);
+  // test a value larger than max(float)
+  assert(__builtin_round(3.4e39 + v) == 3.4e39);
+}
+
+int main(int argc, char* argv[]) {
+  float host_value = 0;
+
+  // Launch the kernel.
+  test_round<<<1, 1>>>(0);
+  cudaError_t err = cudaDeviceSynchronize();
+  if (err != cudaSuccess) {
+    printf("CUDA error %d\n", (int)err);
+    return 1;
+  }
+
+  printf("Success!\n");
+  return 0;
+}