How acos/asin/atan in pytorch handle half input on gpu?

garfieldxs · October 31, 2023, 5:50am

I see that acos for gpu has implementation here.

pytorch/pytorch/blob/08dbfecdbdf2af6f66b3226881c71d8977431197/aten/src/ATen/native/cuda/UnaryGeometricAcosKernel.cu#L50


      
                  });
          #endif
            } else {
              AT_DISPATCH_FLOATING_TYPES_AND2(
                  ScalarType::Half,
                  ScalarType::BFloat16,
                  common_dtype,
                  "acos_cuda",
                  [&]() {
                    gpu_kernel(iter, [] GPU_LAMBDA(scalar_t a) -> scalar_t {
                      return ::acos(a);
                    });
                  });
            }
          }
          
          REGISTER_DISPATCH(acos_stub, &acos_kernel_cuda);
          
          } // namespace at::native

But when I tried to compile the following code

// Type your code here, or load an example.
#include <cuda.h>
#include <cuda_fp16.h>
#include <cuda_runtime.h>

__global__ void asin_kernel(half *input, half *output, int size) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < size)  {
        output[tid] = ::acos(input[tid]);
    }
}

I got error

<source>(9): error: more than one instance of overloaded function "acos" matches the argument list:
            function "std::acos(long double)" (declared at line 91 of /opt/compiler-explorer/gcc-10.2.0/include/c++/10.2.0/cmath)
            function "std::acos(float)" (declared at line 87 of /opt/compiler-explorer/gcc-10.2.0/include/c++/10.2.0/cmath)
            argument types are: (half)
          output[tid] = ::acos(input[tid]);
                        ^

1 error detected in the compilation of "<source>".
Compiler returned: 2

how does pytorch get half input with acos work?