[OpenMP][CUDA] Add missing overload for `remquo(float,float,int*)`

[jdoerfert]

Summary

Reported by Colleen Bertoni <[e-mail]> after running the OvO test
suite: https://github.com/TApplencourt/OvO/

The template overload is still hidden behind an ifdef for OpenMP. In the
future we probably want to remove the ifdef but that requires further
testing.

[jdoerfert]

This should fix the last OvO math test (I think C++11) we fail. While those tests are "simple" they are fairly exhaustive and it's a good sign to pass them ;)

[JonChesterfield]

Change obviously good. Am I right in reading this as all of OvO then passes for trunk, nvptx?

[jdoerfert]

In D89971#2347762, @JonChesterfield wrote:

Change obviously good. Am I right in reading this as all of OvO then passes for trunk, nvptx?

C++11 math tests on one system. These are system dependent to some degree so I would not claim it will always work, but it looks promising. and yes, nvptx, though nothing really cares about that, it just resolves down to nv calls instead of xyz calls at the end of the day.

[tra]

The template overload

function overload.

is still hidden behind an ifdef for OpenMP. In the
future we probably want to remove the ifdef but that requires further
testing.

I don't think it's the case. I've just ran clang++ -x cuda /dev/null --cuda-host-only -dD -E and I do see the definitions for tanh just above your change.

[jdoerfert]

In D89971#2347909, @tra wrote:

The template overload

function overload.

is still hidden behind an ifdef for OpenMP. In the
future we probably want to remove the ifdef but that requires further
testing.

I don't think it's the case. I've just ran clang++ -x cuda /dev/null --cuda-host-only -dD -E and I do see the definitions for tanh just above your change.

My comment was about the template I point out in the inline comment below. Also I was talking about the OpenMP declare variant overloads, not the cuda definitions.
TBH, I don't understand what you want to tell me with the remark about the cuda definitions and tanh. Nothing the OpenMP stuff does should impact cuda and tanh was there already, I'm confused :(

clang/lib/Headers/__clang_cuda_cmath.h

330

The template overload I was talking about is this one.

[tra]

clang/lib/Headers/__clang_cuda_cmath.h

330

Ah. It makes more sense now.

The patch description may need to be rephrased. You've added a function overload, but the description talks about the template which is *not* in the patch.
Something like Added function overload for remquo(float...)... the template function overload for other types... is still behind OPENMP_NVPTX

[jdoerfert]

clang/lib/Headers/__clang_cuda_cmath.h

330

Will do:

Not affected by this patch is the `remquo` template overload which is
still hidden behind an ifdef for OpenMP. In the future we probably want
to remove the ifdef but that requires further testing.

[tra]

Looks like it breaks CUDA: http://lab.llvm.org:8011/#/builders/46/builds/304

[10/873] Building CXX object External/CUDA/CMakeFiles/printf-cuda-11.0-c++14-libstdc++-6.dir/printf.cu.o
FAILED: External/CUDA/CMakeFiles/printf-cuda-11.0-c++14-libstdc++-6.dir/printf.cu.o 
/buildbot/cuda-t4-0/work/None_clang-cuda-t4/clang/bin/clang++  -DNDEBUG  -O3 -DNDEBUG   -w -Werror=date-time -UNDEBUG --cuda-path=/buildbot/cuda-t4-0/work/None_clang-cuda-t4/external/cuda/cuda-11.0 -I/buildbot/cuda-t4-0/work/None_clang-cuda-t4/external/cuda/cuda-11.0/include --cuda-gpu-arch=sm_75 -std=c++14 -stdlib=libstdc++ -gcc-toolchain /buildbot/cuda-t4-0/work/None_clang-cuda-t4/external/cuda/gcc-6 -DSTDLIB_VERSION=2014 -MD -MT External/CUDA/CMakeFiles/printf-cuda-11.0-c++14-libstdc++-6.dir/printf.cu.o -MF External/CUDA/CMakeFiles/printf-cuda-11.0-c++14-libstdc++-6.dir/printf.cu.o.d -o External/CUDA/CMakeFiles/printf-cuda-11.0-c++14-libstdc++-6.dir/printf.cu.o -c /buildbot/cuda-t4-0/work/None_clang-cuda-t4/llvm-test-suite/External/CUDA/printf.cu
In file included from <built-in>:1:
In file included from /buildbot/cuda-t4-0/work/None_clang-cuda-t4/clang/lib/clang/12.0.0/include/__clang_cuda_runtime_wrapper.h:412:
/buildbot/cuda-t4-0/work/None_clang-cuda-t4/clang/lib/clang/12.0.0/include/__clang_cuda_cmath.h:177:18: error: redefinition of 'remquo'
__DEVICE__ float remquo(float __n, float __d, int *__q) {
                 ^
/buildbot/cuda-t4-0/work/None_clang-cuda-t4/external/cuda/cuda-11.0/include/crt/math_functions.hpp:685:31: note: previous definition is here
__MATH_FUNCTIONS_DECL__ float remquo(float a, float b, int *quo)
                              ^
1 error generated when compiling for sm_75.
[11/873] Building CXX object External/CUDA/CMakeFiles/builtin_var-cuda-11.0-c++14-libstdc++-7.dir/builtin_var.cu.o
FAILED: External/CUDA/CMakeFiles/builtin_var-cuda-11.0-c++14-libstdc++-7.dir/builtin_var.cu.o 
/buildbot/cuda-t4-0/work/None_clang-cuda-t4/clang/bin/clang++  -DNDEBUG  -O3 -DNDEBUG   -w -Werror=date-time -UNDEBUG --cuda-path=/buildbot/cuda-t4-0/work/None_clang-cuda-t4/external/cuda/cuda-11.0 -I/buildbot/cuda-t4-0/work/None_clang-cuda-t4/external/cuda/cuda-11.0/include --cuda-gpu-arch=sm_75 -std=c++14 -stdlib=libstdc++ -gcc-toolchain /buildbot/cuda-t4-0/work/None_clang-cuda-t4/external/cuda/gcc-7 -DSTDLIB_VERSION=2014 -MD -MT External/CUDA/CMakeFiles/builtin_var-cuda-11.0-c++14-libstdc++-7.dir/builtin_var.cu.o -MF External/CUDA/CMakeFiles/builtin_var-cuda-11.0-c++14-libstdc++-7.dir/builtin_var.cu.o.d -o External/CUDA/CMakeFiles/builtin_var-cuda-11.0-c++14-libstdc++-7.dir/builtin_var.cu.o -c /buildbot/cuda-t4-0/work/None_clang-cuda-t4/llvm-test-suite/External/CUDA/builtin_var.cu
In file included from <built-in>:1:
In file included from /buildbot/cuda-t4-0/work/None_clang-cuda-t4/clang/lib/clang/12.0.0/include/__clang_cuda_runtime_wrapper.h:412:
/buildbot/cuda-t4-0/work/None_clang-cuda-t4/clang/lib/clang/12.0.0/include/__clang_cuda_cmath.h:177:18: error: redefinition of 'remquo'
__DEVICE__ float remquo(float __n, float __d, int *__q) {
                 ^
/buildbot/cuda-t4-0/work/None_clang-cuda-t4/external/cuda/cuda-11.0/include/crt/math_functions.hpp:685:31: note: previous definition is here
__MATH_FUNCTIONS_DECL__ float remquo(float a, float b, int *quo)
                              ^
1 error generated when compiling for sm_75.
[12/873] Building CXX object External/CUDA/CMakeFiles/assert-cuda-11.0-c++14-libstdc++-7.dir/assert.cu.o
FAILED: External/CUDA/CMakeFiles/assert-cuda-11.0-c++14-libstdc++-7.dir/assert.cu.o 
/buildbot/cuda-t4-0/work/None_clang-cuda-t4/clang/bin/clang++  -DNDEBUG  -O3 -DNDEBUG   -w -Werror=date-time -UNDEBUG --cuda-path=/buildbot/cuda-t4-0/work/None_clang-cuda-t4/external/cuda/cuda-11.0 -I/buildbot/cuda-t4-0/work/None_clang-cuda-t4/external/cuda/cuda-11.0/include --cuda-gpu-arch=sm_75 -std=c++14 -stdlib=libstdc++ -gcc-toolchain /buildbot/cuda-t4-0/work/None_clang-cuda-t4/external/cuda/gcc-7 -DSTDLIB_VERSION=2014 -MD -MT External/CUDA/CMakeFiles/assert-cuda-11.0-c++14-libstdc++-7.dir/assert.cu.o -MF External/CUDA/CMakeFiles/assert-cuda-11.0-c++14-libstdc++-7.dir/assert.cu.o.d -o External/CUDA/CMakeFiles/assert-cuda-11.0-c++14-libstdc++-7.dir/assert.cu.o -c /buildbot/cuda-t4-0/work/None_clang-cuda-t4/llvm-test-suite/External/CUDA/assert.cu
In file included from <built-in>:1:
In file included from /buildbot/cuda-t4-0/work/None_clang-cuda-t4/clang/lib/clang/12.0.0/include/__clang_cuda_runtime_wrapper.h:412:
/buildbot/cuda-t4-0/work/None_clang-cuda-t4/clang/lib/clang/12.0.0/include/__clang_cuda_cmath.h:177:18: error: redefinition of 'remquo'
__DEVICE__ float remquo(float __n, float __d, int *__q) {
                 ^
/buildbot/cuda-t4-0/work/None_clang-cuda-t4/external/cuda/cuda-11.0/include/crt/math_functions.hpp:685:31: note: previous definition is here
__MATH_FUNCTIONS_DECL__ float remquo(float a, float b, int *quo)
                              ^
1 error generated when compiling for sm_75.

[tra]

Never mind. I see that you've fixed it in 17c8251bca83072d2f3e00f936d6ce24500e6b02

[jdoerfert]

In D89971#2359398, @tra wrote:

Never mind. I see that you've fixed it in 17c8251bca83072d2f3e00f936d6ce24500e6b02

Yeah, as I write in the commit message, this confused me but the short term solution was simple so I went ahead.

[tra]

I'll take a look at what broke the original patch.

clang/lib/Headers/__clang_cuda_cmath.h

 /*===---- __clang_cuda_cmath.h - Device-side CUDA cmath support ------------===
  *
  * 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
  *
  *===-----------------------------------------------------------------------===
  */
 #ifndef __CLANG_CUDA_CMATH_H__
 #define __CLANG_CUDA_CMATH_H__
 #ifndef __CUDA__
 #error "This file is for CUDA compilation only."
 #endif
 
 #ifndef __OPENMP_NVPTX__
 #include <limits>
 #endif
 
 // CUDA lets us use various std math functions on the device side.  This file
 // works in concert with __clang_cuda_math_forward_declares.h to make this work.

 // std covers all of the known knowns.
 
 #ifdef __OPENMP_NVPTX__
 #define __DEVICE__ static constexpr __attribute__((always_inline, nothrow))
 #else
 #define __DEVICE__ static __device__ __inline__ __attribute__((always_inline))
 #endif
 
 __DEVICE__ long long abs(long long __n) { return ::llabs(__n); }
 __DEVICE__ long abs(long __n) { return ::labs(__n); }
 __DEVICE__ float abs(float __x) { return ::fabsf(__x); }
 __DEVICE__ double abs(double __x) { return ::fabs(__x); }
 __DEVICE__ float acos(float __x) { return ::acosf(__x); }
 __DEVICE__ float asin(float __x) { return ::asinf(__x); }
 __DEVICE__ float atan(float __x) { return ::atanf(__x); }
 __DEVICE__ float atan2(float __x, float __y) { return ::atan2f(__x, __y); }
 __DEVICE__ float ceil(float __x) { return ::ceilf(__x); }
 __DEVICE__ float cos(float __x) { return ::cosf(__x); }
 __DEVICE__ float cosh(float __x) { return ::coshf(__x); }
 __DEVICE__ float exp(float __x) { return ::expf(__x); }
 __DEVICE__ float fabs(float __x) { return ::fabsf(__x); }
 __DEVICE__ float floor(float __x) { return ::floorf(__x); }
 __DEVICE__ float fmod(float __x, float __y) { return ::fmodf(__x, __y); }
 __DEVICE__ int fpclassify(float __x) {
   return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,

 }
 __DEVICE__ bool signbit(float __x) { return ::__signbitf(__x); }
 __DEVICE__ bool signbit(double __x) { return ::__signbitd(__x); }
 __DEVICE__ float sin(float __x) { return ::sinf(__x); }
 __DEVICE__ float sinh(float __x) { return ::sinhf(__x); }
 __DEVICE__ float sqrt(float __x) { return ::sqrtf(__x); }
 __DEVICE__ float tan(float __x) { return ::tanf(__x); }
 __DEVICE__ float tanh(float __x) { return ::tanhf(__x); }
+__DEVICE__ float remquo(float __n, float __d, int *__q) {
+  return ::remquof(__n, __d, __q);
+}
 
 // Notably missing above is nexttoward.  We omit it because
 // libdevice doesn't provide an implementation, and we don't want to be in the
 // business of implementing tricky libm functions in this header.
 
 #ifndef __OPENMP_NVPTX__
 
 // Now we've defined everything we promised we'd define in
ldexp(__T __x, int __exp) {
   return std::ldexp((double)__x, __exp);
 }
 
 template <typename __T1, typename __T2>
 __DEVICE__ typename __clang_cuda_enable_if<
     std::numeric_limits<__T1>::is_specialized &&
         std::numeric_limits<__T2>::is_specialized,
     double>::type
 remquo(__T1 __x, __T2 __y, int *__quo) {
   return std::remquo((double)__x, (double)__y, __quo);
 }
 
 template <typename __T>
 __DEVICE__ typename __clang_cuda_enable_if<std::numeric_limits<__T>::is_integer,
                                            double>::type
 scalbln(__T __x, long __exp) {
   return std::scalbln((double)__x, __exp);