Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

OpenCL compiler does not support fract() #101

Open
gerd-sk opened this issue Oct 31, 2022 · 2 comments
Open

OpenCL compiler does not support fract() #101

gerd-sk opened this issue Oct 31, 2022 · 2 comments
Labels

Comments

@gerd-sk
Copy link

gerd-sk commented Oct 31, 2022

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 AmitBM added the bug label Dec 8, 2022
@AmitBM
Copy link
Contributor

AmitBM commented Dec 15, 2022

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

Can you please give it a shot and confirm?

@gerd-sk
Copy link
Author

gerd-sk commented Dec 15, 2022

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.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
Projects
None yet
Development

No branches or pull requests

2 participants