[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

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

17

__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

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

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 ↗

(On Diff #192762)

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(round 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/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.
+
+template <typename T>
+__global__ void round(T* x, T* y) {
+  y[threadIdx.x] = __builtin_roundf(x[threadIdx.x]);
+}
+
+const int kDataLen = 4;
+
+template <typename T>
+void test_round(T* host_x, T* expected_y) {
+  T host_y[kDataLen];
+
+  // Copy input data to device.
+  T* device_x;
+  T* device_y;
+  cudaMalloc(&device_x, kDataLen * sizeof(T));
+  cudaMalloc(&device_y, kDataLen * sizeof(T));
+  cudaMemcpy(device_x, host_x, kDataLen * sizeof(T),
+             cudaMemcpyHostToDevice);
+
+  // Launch the kernel.
+  round<T><<<1, kDataLen>>>(device_x, device_y);
+
+  // Copy output data to host.
+  cudaDeviceSynchronize();
+  cudaMemcpy(host_y, device_y, kDataLen * sizeof(T),
+             cudaMemcpyDeviceToHost);
+
+  // Print and compare the results.
+  for (int i = 0; i < kDataLen; ++i) {
+    std::cout << "y[" << i << "] = " << host_y[i] << " " << expected_y[i]
+      << " " << host_y[i] - expected_y[i] << "\n";
+    assert(abs((host_y[i] - expected_y[i])/expected_y[i]) < 0.001);
+  }
+}
+
+int main(int argc, char* argv[]) {
+
+  float float_x[kDataLen] = {-0.5f, 8.5f, -8.38861e+06f, 8.38861e+06f};
+  float float_y[kDataLen] = {-1.0f, 9.0f, -8.38861e+06f, 8.38861e+06f};
+  test_round(float_x, float_y);
+
+  double double_x[kDataLen] = {0.5, -8.5, 4.5035996e+15, -4.5035996e+15};
+  double double_y[kDataLen] = {1.0, -9.0, 4.5035996e+15, -4.5035996e+15};
+  test_round(double_x, double_y);
+
+  cudaDeviceReset();
+  return 0;
+}