[libspirv] Define schar overloads via remangling; not source

[frasercrmck]

We were previously achieving the signed char builtin definitions in libspirv via one of two ways. The first was explicitly definining schar overloads of builtins in the source. The second was by remangling 'char' builtins to one of schar or uchar, depending on the host platform.

Since we are defining our builtins in OpenCL C, the plain 'char' type is already a signed type. This presents us with the opportunity to achieve our desired schar builtins solely through remangling. The primary idea is to reduce our libclc/libspirv diff with upstream. We also have the option to introduce signed char builtins upstream. As it stands the schar problem isn't far from the 'half' mangling problem that we also now deal with purely in the remangler.

[frasercrmck]

CC @wenju-he

[ldrumm]

Since we are defining our builtins in OpenCL C, the plain 'char' type is already a signed type.

Does that include ARM?

[frasercrmck]

Since we are defining our builtins in OpenCL C, the plain 'char' type is already a signed type.

Does that include ARM?

To be honest I'm not sure. It should be given the OpenCL C specification (here). I'd say that if ARM is treating char as unsigned for OpenCL code it's a clang (or a user) bug.

[ldrumm]

I'd say that if ARM is treating char as unsigned for OpenCL code it's a clang (or a user) bug.

Agreed. However, clang OpenCL has historically been maintained primarily by ARM, and I think we encountered this before in ComputeAorta. I'm not saying it's a blocker, but it may be worth checking

[ldrumm]

I'd say that if ARM is treating char as unsigned for OpenCL code it's a clang (or a user) bug.

Agreed. However, clang OpenCL has historically been maintained primarily by ARM, and I think we encountered this before in ComputeAorta. I'm not saying it's a blocker, but it may be worth checking

Since we are defining our builtins in OpenCL C, the plain 'char' type is already a signed type.

Does that include ARM?

Resolved via offline demonstration. LGTM

[frasercrmck]

I'd say that if ARM is treating char as unsigned for OpenCL code it's a clang (or a user) bug.

Agreed. However, clang OpenCL has historically been maintained primarily by ARM, and I think we encountered this before in ComputeAorta. I'm not saying it's a blocker, but it may be worth checking

(As discussed offline)

Compiling a simple C and OpenCL C file which casts a char to int shows that OpenCL does indeed seem to be guaranteeing signed chars by default, even for targets like arm and riscv64 where the C code is unsigned.

[wenju-he]

thanks @frasercrmck. This is very nice.

[frasercrmck]

ping, thanks.

[Naghasan]

LGTM, glad you addressed this legacy code

[frasercrmck]

ping @intel/dpcpp-cfe-reviewers, thanks

[hvdijk]

To be honest I'm not sure. It should be given the OpenCL C specification

Whether char is signed in SYCL code is independent of whether it is signed in OpenCL C code. SYCL has whatever the host platform has. But it seems like this only affects functions that don't care whether char is signed, in which case that part is fine, just for a different reason.

But something seems to be resulting in a mismatch between what Clang emits and what libclc defines, at least for NativeCPU: we are now getting errors such as

ld.lld: error: undefined symbol: __spirv_GroupAsyncCopy(int, signed char vector[16] AS1*, signed char vector[16] const AS3*, unsigned long, unsigned long, ocl_event)
>>> referenced by llvm-link
>>>               /tmp/test_nd_item.aarch64.-798124.img:(typeinfo name for nd_item_async_work_group_copy::kernel_type<sycl::_V1::vec<signed char, 16>, 1> (.NativeCPUKernel))
>>> referenced by llvm-link
>>>               /tmp/test_nd_item.aarch64.-798124.img:(typeinfo name for nd_item_async_work_group_copy::kernel_type<sycl::_V1::vec<signed char, 16>, 1> (.NativeCPUKernel))
>>> referenced by llvm-link
>>>               /tmp/test_nd_item.aarch64.-798124.img:(typeinfo name for nd_item_async_work_group_copy::kernel_type<sycl::_V1::vec<signed char, 16>, 2> (.NativeCPUKernel))
>>> referenced 3 more times

when building SYCL-CTS for NativeCPU. Am I understanding correctly that the idea is what libclc only explicitly defines plain char versions of these functions, and the remangler is responsible for creating a signed char version from this? If so, that does not appear to be happening.

[frasercrmck]

when building SYCL-CTS for NativeCPU. Am I understanding correctly that the idea is what libclc only explicitly defines plain char versions of these functions, and the remangler is responsible for creating a signed char version from this?

Yes, that's the idea, thanks for reporting the issue.

I'll take a look. It might be that the SPIRVBuiltins.td still needs to offer signed char variants.

[frasercrmck]

Hmm yes, looking at __spirv_GroupAsyncCopy it is using AGenTypeN which includes TLAll which was defined as def TLAll : TypeList<[Char, UChar, Short, (...).

I didn't change this list. It was always using Char which (confusingly and I think incorrectly) was defined as SignedCharTy. When I changed Char to be a proper char I think I inadvertently removed the declarations of all builtins for when you use signed char explicitly. Unfortunately, it seems, there were no LIT tests to check this?

I hope that adding back SChar and adding it to all the all/gentype lists we can regain that functionality.

[frasercrmck]

@hvdijk hopefully #18807 fixes this issue.

[hvdijk]

Thanks! I'm running this on our internal CI and will comment on the PR based on what results I get.