radeon_gpu_analyzer icon indicating copy to clipboard operation
radeon_gpu_analyzer copied to clipboard

OpenCL compiler does not support fract()

Open gerd-sk opened this issue 3 years ago • 2 comments

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.

gerd-sk avatar Oct 31 '22 15:10 gerd-sk

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

Can you please give it a shot and confirm?

AmitBM avatar Dec 15 '22 15:12 AmitBM

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.

gerd-sk avatar Dec 15 '22 17:12 gerd-sk