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

Artem Belevich via Phabricator via llvm-commits llvm-commits at lists.llvm.org
Thu Mar 28 15:22:29 PDT 2019


tra added inline comments.


================
Comment at: External/CUDA/round.cu:16
+template <typename T>
+__global__ void round(T* x, T* y) {
+  y[threadIdx.x] = __builtin_roundf(x[threadIdx.x]);
----------------
I'd use a name different from the standard libm function. `do_round` perhaps?


================
Comment at: External/CUDA/round.cu:17
+__global__ void round(T* x, T* y) {
+  y[threadIdx.x] = __builtin_roundf(x[threadIdx.x]);
+}
----------------
`__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()


================
Comment at: External/CUDA/round.cu:24-47
+  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));
----------------
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. 



================
Comment at: External/CUDA/round.cu:52-53
+
+  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);
----------------
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.


Repository:
  rT test-suite

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D59950/new/

https://reviews.llvm.org/D59950





More information about the llvm-commits mailing list