[CIR][CIRGen] Fix builtin IIT Integer signedness using AST information

[RiverDave]

I found a blocker when working on PTX intrinsics.

One of the problems we have whenever we want to emit generic builtins with the current infrastructure (IIT Descriptors) is that it doesn't allow us to determine signedness in scalar types. When we emit these builtins we seem to type check param/return types:
if (argType != intrinsicType.getInput(i)). In our current approach to decode a type from IIT we force signedness:

static mlir::Type
decodeFixedType(ArrayRef<llvm::Intrinsic::IITDescriptor> &infos,
                mlir::MLIRContext *context) {
  using namespace llvm::Intrinsic;

  IITDescriptor descriptor = infos.front();
  infos = infos.slice(1);

  switch (descriptor.Kind) {
...
  case IITDescriptor::Integer:
    return IntType::get(context, descriptor.Integer_Width, /*isSigned=*/true);

I was finding cases where the function definition would specify x parameter as unsigned but when calling this function with the right type, it would trigger an NYI statement we had encoded(This is not case for OG as by design, LLVMIR does not contain that information):

if (argType != intrinsicType.getInput(i))
        llvm_unreachable("NYI");

Example issue:

  def __nvvm_bar_warp_sync : NVPTXBuiltinPTX<"void(unsigned int)", PTX60>;

__device__ void nvvm_sync(unsigned int mask) {
    __nvvm_bar_warp_sync(mask);  // expects unsigned, but IIT returns signed (wrong assumption)
}

To solve this, I've used a very naive approach, we know that front-end contains all signedness information we need. I've used AST information (I've specifically extracted the Function Decl) to determine whether if we should enforce signedness or not based on the definition it was declared with.

I've also added the respective CUDA test that was being affected by this issue. along with the implementing the rest of the NYI block related to this.

[github-actions]

✅ With the latest revision this PR passed the C/C++ code formatter.

[bcardosolopes]

Ok, LGTM once a few final nits get address - thanks for your patience