OpenCL compiler does not support fract()

[gerd-sk]

If I try to compile an OpenCL kernel that calls the fract() function, I get the following error:

ld.lld: error: undefined hidden symbol: fract(float, float AS0*)

referenced by fract.cl:7 (/home/EU/despalig/Documents/RadeonGPUAnalyzer/projects/221026-150246/Clone0/fract.cl:7)
/tmp/despalig/fract-220ba7.o:(src_MyKernel)
referenced by fract.cl:7 (/home/EU/despalig/Documents/RadeonGPUAnalyzer/projects/221026-150246/Clone0/fract.cl:7)
/tmp/despalig/fract-220ba7.o:(src_MyKernel)
fract.cl.txt

The kernel code is:
__kernel void src_MyKernel(__global float *out, __global const float *in)
{
int index = get_global_linear_id();
float dummy;

out[index] = fract(in[index], &dummy);

}

I have used RGA shipped with RadeonDeveloperToolSuite-2022-08-01-115. Application version: 2.6.2.38 (07/25/2022).
I am running this on Ubuntu 20.04 Linux.

This kernel runs fine on my machine. Offline analysis with RGA fails. I suspect the OpenCL libraries shipped with RGA are obsolete.

[AmitBM]

This issue should be resolved in RGA 2.7 which was released yesterday.

Can you please give it a shot and confirm?

[gerd-sk]

RGA 2.7 resolves the fract() issue. Thank you for fixing the problem.

But now the analysis result is much different from what I got with RGA 2.6.2.38
For a larger kernel, with asic gfx906 I got:
VGPRs: 63 / 256 | SGPRs: 50 / 102 with RGA 2.6.2.38
VGPRs: 83 / 256 | SGPRs: 48 / 102 with RGA 2.7.0.117
The build settings were identical in both versions (basically all default or empty except for the target GPU, which I set to gfx906)
And the code size of the binary became larger with the new version.
Honestly speaking, compilation results of RGA 2.6.2.38 look better to me.

[AmitBM]

Hi @gerd-sk, if this is still an issue, please create a ticket for the ROCm compiler team here.