[HIP] Add generic address space impl of __spirv_ocl_frexp for HIP backend

[FMarno]

This involved manually mangling the namespaces since the default would be address space 5 otherwise.

[bader]

LGTM, although vector variant implementations are verbose and it would be great to re-use vector functions instead of scalar versions.

[bader]

Can't we use VEC_FUNCNAME(2) here? Similar to the existing approach where computation for vector length 2*N done by combining results for vector length N.

[FMarno]

Yes, that is possible but would it not lead to extra vector being created? For example VEC_FUNCNAME(16) would create a 8 int2, 4 int4, 2 int8 and an int16.

[bader]

VEC_FUNCNAME(16) should create just 2 int8. Currently it creates 16(!) scalar values, so I don't think creating 2 int8 should be a problem.

[bader]

Sorry, I think I didn't get your concern right. There should be no overhead on creating vectors. Compiler should generate the same for code for both variants. It's just less code to read.

[FMarno]

Ok, looking at _CLC_V_V_VP_VECTORIZE and modf this also seems to be a fairly common pattern. Thanks!

[FMarno]

Apologies, I accidentally rebased. It's almost all redone anyway.
@bader I added a new file with a common macro but I'm not sure if I've put it in the correct place or if this needs to be added to the SOURCES file or CMake somewhere, can you please advise.

[bader]

I'm not sure if I've put it in the correct place or if this needs to be added to the SOURCES file or CMake somewhere, can you please advise.

I suggest we change the file extension from .cl to .h for mangle_common.cl.

According to my understanding SOURCES is used to enable the library binary re-build when sources are changed, so adding new files impacting re-build results is required.
Everything else looks good.

[bader]

Suggested change

	#include "mangle_common.cl"
	#include "mangle_common.h"

[FMarno]

I think the problem is that modf and sincos are not inlining. One thing I've notices is that the test seems to have a slightly different ending, "S_" instead of "Dv2_f", which seems to be something to do with compression in Itanium (https://itanium-cxx-abi.github.io/cxx-abi/abi.html#mangling-compression), caused by the type float2 being using twice. That would also explain why frexp doesn't have the problem, since it have two different argument types. I'll try fixing it in the morning 😄

[FMarno]

Hi @bader,
After my change I reran the test on my local machine and it seems to be inlining now, which is good.

I noticed that modf, sincos and frexp dont work for the half type, but this isn't just limited to the HIP backend.
minimal reproducer:

#include <CL/sycl.hpp>

int main() {
    sycl::half x;
    sycl::modf(sycl::half{1.0}, &x);
    sycl::sincos(sycl::half{1.0}, &x);
    return 0;
}

Is it ok if open an issue for this?

[bader]

Is it ok if open an issue for this?

Yes. Please, file a bug report.

[FMarno]

Bug reported here: #5418