Implement operators logical_and and logical_or

[xuhdev]

Stack from ghstack:

• Implement operators logical_and and logical_or #24379 Implement operators logical_and and logical_or

Pull Request resolved: #24379

Differential Revision: D16830128

[xuhdev]

@bddppq Are you aware of a nvfunctional equivalent in ROCm?

[gchanan]

I don't think we should do this; not only does it not work with ROCm, but it can't be inlined.

[xuhdev]

I now changed this to templating. According to my experiments, nvcc seems to be able to inline functions passed in as template parameters.

__host__ __device__ float called1(float x) {
  return x + 1;
}

__host__ __device__ float called2(float x) {
  return x + 10;
}

template <typename Op>
__host__ __device__ float caller(float x, Op op) {
  return op(x);
}

__global__ void user(float x, float& y) {
  x = caller(x, called1);
  y = caller(x, called2);
}

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

//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-24817639
// Cuda compilation tools, release 10.0, V10.0.130
// Based on LLVM 3.4svn
//

.version 6.3
.target sm_60
.address_size 64

	// .globl	_Z4userfRf

.visible .entry _Z4userfRf(
	.param .f32 _Z4userfRf_param_0,
	.param .u64 _Z4userfRf_param_1
)
{
	.reg .f32 	%f<4>;
	.reg .b64 	%rd<3>;


	ld.param.f32 	%f1, [_Z4userfRf_param_0];
	ld.param.u64 	%rd1, [_Z4userfRf_param_1];
	cvta.to.global.u64 	%rd2, %rd1;
	add.f32 	%f2, %f1, 0f3F800000;
	add.f32 	%f3, %f2, 0f41200000;
	st.global.f32 	[%rd2], %f3;
	ret;
}