Move the CUDA implementation of trunc to ATen.

[xuhdev]

Stack from ghstack:

• Move the CUDA implementation of log to ATen. #26494 Move the CUDA implementation of log to ATen.
• Simplify operator sign using the helper. #25592 Simplify operator sign using the helper.
• Move the CUDA implementation of trunc to ATen. #25423 Move the CUDA implementation of trunc to ATen.

Fix #24650

Differential Revision: D17397489

[VitalyFedyunin]

Are you sure there is no need to overload to call truncf ?

[xuhdev]

Yeah, it's overloaded https://en.cppreference.com/w/cpp/numeric/math/trunc (also search for VSTD::trunc in crt/math_functions.h in the cuda include dir)

[xuhdev]

If you are worried, here's some evidence. Compile the following program using

nvcc -ptx -src-in-ptx -arch=sm_60 test.cu

#include <cuda_runtime.h>

__global__ void test_trunc_f(float x,  float& x2) {
  x2 = std::trunc(x);
}

__global__ void test_trunc_d(double x,  double& x2) {
  x2 = std::trunc(x);
}

__global__ void test_trunc_f_d(float x,  float& x2) {
  x2 = truncf(x);
}

The output shows (Note that the first and third functions are compiled to the same asm code)

//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-26907403
// Cuda compilation tools, release 10.1, V10.1.243
// Based on LLVM 3.4svn
//

.version 6.4
.target sm_60
.address_size 64

	// .globl	_Z12test_trunc_ffRf

.visible .entry _Z12test_trunc_ffRf(
	.param .f32 _Z12test_trunc_ffRf_param_0,
	.param .u64 _Z12test_trunc_ffRf_param_1
)
{
	.reg .f32 	%f<3>;
	.reg .b64 	%rd<3>;


	ld.param.f32 	%f1, [_Z12test_trunc_ffRf_param_0];
	ld.param.u64 	%rd1, [_Z12test_trunc_ffRf_param_1];
	cvta.to.global.u64 	%rd2, %rd1;
	cvt.rzi.f32.f32	%f2, %f1;
	st.global.f32 	[%rd2], %f2;
	ret;
}

	// .globl	_Z12test_trunc_ddRd
.visible .entry _Z12test_trunc_ddRd(
	.param .f64 _Z12test_trunc_ddRd_param_0,
	.param .u64 _Z12test_trunc_ddRd_param_1
)
{
	.reg .f64 	%fd<3>;
	.reg .b64 	%rd<3>;


	ld.param.f64 	%fd1, [_Z12test_trunc_ddRd_param_0];
	ld.param.u64 	%rd1, [_Z12test_trunc_ddRd_param_1];
	cvta.to.global.u64 	%rd2, %rd1;
	cvt.rzi.f64.f64	%fd2, %fd1;
	st.global.f64 	[%rd2], %fd2;
	ret;
}

	// .globl	_Z14test_trunc_f_dfRf
.visible .entry _Z14test_trunc_f_dfRf(
	.param .f32 _Z14test_trunc_f_dfRf_param_0,
	.param .u64 _Z14test_trunc_f_dfRf_param_1
)
{
	.reg .f32 	%f<3>;
	.reg .b64 	%rd<3>;


	ld.param.f32 	%f1, [_Z14test_trunc_f_dfRf_param_0];
	ld.param.u64 	%rd1, [_Z14test_trunc_f_dfRf_param_1];
	cvta.to.global.u64 	%rd2, %rd1;
	cvt.rzi.f32.f32	%f2, %f1;
	st.global.f32 	[%rd2], %f2;
	ret;
}

rocm failure looks reasonable

[VitalyFedyunin]

Please fix pr/py2-clang7-rocmdeb-ubuntu16.04

[xuhdev]

@VitalyFedyunin Updated. Should have fixed it now

[VitalyFedyunin]

"All checks have passed" super suspicious ;)

[xuhdev]

@VitalyFedyunin A friendly reminder that the previous two merged commits in this stack seem to have put the incorrect authorship: f55a9da

[VitalyFedyunin]

Please rebase ( contention on the UnaryOps.cpp is way to high )

[xuhdev]

Done!

[facebook-github-bot]

@VitalyFedyunin merged this pull request in 7bdc0c1.

[xuhdev]

Thanks!