[CIR] Fix handling of varargs and void-return builtins in intrinsic t…

[mahmood82]

Fix several issues in CIR intrinsic type construction for built-in functions such as printf (variadic) and void-return intrinsics.

• Treat IITDescriptor::VarArg as a terminator marker rather than a type. It is now handled correctly in getIntrinsicType instead of decodeFixedType.
• Added support for marking CIR function types as variadic.
• Added proper handling of builtins returning void: CIRFuncType now omits an explicit void return type when appropriate.
• Ensures correct CIR signatures for OpenCL and C built-ins that use ellipsis (printf) or return void (e.g. printf).

[bcardosolopes]

Great! Also needs a testcase

[mahmood82]

Unfortunately I was not able to create a unit test that fails.

I made a simple printf("...") test and it fails on the local repository of our project. Because we have a private version of print if. The CIR lowering looks like this:

%6 = cir.llvm.intrinsic "private.printf" %5 : (!cir.ptr<!s8i>) -> !void

However using the original incubator repository the default built-in printf was used, so I did not hit the new code in this PR:

%5 = cir.call @printf(%4) : (!cir.ptr<!s8i, addrspace(offload_constant)>) -> !s32i cc(spir_function) extra(#fn_attr)

The change I made in this PR indeed fixes the private intrinsic case.
P.S. probably this change (or part of it) relates to #1685.

[bcardosolopes]

However using the original incubator repository the default built-in printf was used

Is this right? Is this what OG does? Looks like you need to fix more things in this PR so it can introduce a testcase?

[mahmood82]

However using the original incubator repository the default built-in printf was used

Is this right? Is this what OG does? Looks like you need to fix more things in this PR so it can introduce a testcase?

The issue is that getIntrinsicType in CIRGenBuiltin.cpp is only exercised for LLVM intrinsics, not normal builtins like printf. In our private repo, __private_printf is defined as a custom intrinsic, which triggers the bug.
Upstream Clang has no equivalent intrinsic for printf, so a normal printf test won't hit this code. To test it upstream, we'd need a kernel that calls an LLVM intrinsic.
The difference is due to lowering, not the correctness of the PR.

[bcardosolopes]

The difference is due to lowering, not the correctness of the PR.

Gotcha, thanks for the explanation

To test it upstream, we'd need a kernel that calls an LLVM intrinsic.

Can't you synthetize that with a LLVM intrinsic that can be used within a kernel? (assuming one exists?)

[mahmood82]

The difference is due to lowering, not the correctness of the PR.

Gotcha, thanks for the explanation

To test it upstream, we'd need a kernel that calls an LLVM intrinsic.

Can't you synthetize that with a LLVM intrinsic that can be used within a kernel? (assuming one exists?)

In this case printf is unique, I did not find other variadic built-in. I tried this but unfortunately didn't hit our flow:

extern void __builtin_myvararg(void *format, ...);
__kernel void test() {
  __builtin_myvararg(0);
}