[BUG]: Invalid instruction generated by `cuda::ptx::shr` for non-stdint type

[davebayer]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct

Type of Bug

Silent Failure

Component

libcu++

Describe the bug

cuda::ptx::shr provides generic overloads for binary types and then provides specializations for cuda::std::int(16, 32, 64)_t types that result in an arithmetic shift instead of a binary shift.

The problem is that stdint types are aliases to standard integer types and int64_t can be either long (Linux) or long long (Windows), which leads to us not providing an overload for the second type.

That means that if I'm on a linux machine and I pass long long to cuda::ptx::shr, I get binary shift, which is unexpected.

How to Reproduce

A simple program like:

#include <cuda/ptx>

__device__ auto fn(long v, unsigned s)
{
    return cuda::ptx::shr(v, s);
}

__device__ auto fn(long long v, unsigned s)
{
    return cuda::ptx::shr(v, s);
}

result in:

.visible .func  (.param .b64 func_retval0) _Z2fnlj(
	.param .b64 _Z2fnlj_param_0,
	.param .b32 _Z2fnlj_param_1
)
{
	.reg .b32 	%r<2>;
	.reg .b64 	%rd<3>;


	ld.param.u64 	%rd2, [_Z2fnlj_param_0];
	ld.param.u32 	%r1, [_Z2fnlj_param_1];
	// begin inline asm
	shr.s64 %rd1, %rd2, %r1;
	// end inline asm
	st.param.b64 	[func_retval0+0], %rd1;
	ret;

}
	// .globl	_Z2fnxj
.visible .func  (.param .b64 func_retval0) _Z2fnxj(
	.param .b64 _Z2fnxj_param_0,
	.param .b32 _Z2fnxj_param_1
)
{
	.reg .b32 	%r<2>;
	.reg .b64 	%rd<3>;


	ld.param.u64 	%rd2, [_Z2fnxj_param_0];
	ld.param.u32 	%r1, [_Z2fnxj_param_1];
	// begin inline asm
	shr.b64 %rd1, %rd2, %r1;                        // <------- should be .s64
	// end inline asm
	st.param.b64 	[func_retval0+0], %rd1;
	ret;

}

Expected behavior

All signed integer types should do arithmetic shift.

Reproduction link

https://godbolt.org/z/zWP3PK77E

Operating System

No response

nvidia-smi output

No response

NVCC version

No response