AnsweredAssumed Answered

Using AMD_IL code under OpenCL runtime environment.

Question asked by realhet on Jan 6, 2013
Latest reply on Jan 14, 2013 by coordz

Hi,

 

In order to make my stuff up to date, I'm leaving CAL behind and trying to use OpenCL for buffer/kernel management.

My goal is to execute an amd_il source code under the OpenCL api, so I can still code in low level and also access all the older cards with a WORKING api.

 

In order to reach the goal, I did the following:

1. compile AMD_IL source with calcl, make a cal_elf image. (Thats the only part I somewhat rely on good old cal)

2. compile a skeleton opencl source to an ocl elf image. (skeleton code 'defines' the buffers (uav,cb))

3. replace the cal_elf section inside the ocl_elf with the generated cal_elf image (step 1).

 

At the moment, the skeleton code is as simple as this:

__kernel void main(__global int* uav,__constant int* cb) { uav[get_global_id(0)]+=cb[0]; }

//one uav, one cb, and a global id, that's all.

//this code is modifyes(!) the contents of the uav

 

It consists of 3 basic operations: cb.read, uav.read, uav.write

 

I took out the .il source of this small opencl program, and recompiled with calcl. Inserted to the skeleton ocl_elf, and gave it to OpenCL API, the results was like:

All operations worked well except the uav.read which is always returned zero.

 

I investigated further and found a difference between the IL->ISA compilers of OpenCL and CAL. (I can only use explicitly CAL's compiler)

This difference was the VFETCH instruction's BUFFER_ID value:

calCl: VFETCH R0.x___, R1.x, fc0

OpenCL's internal compiler: VFETCH R0.x___, R1.x, fc153

(note that it was the result of the same .il program)

 

I don't want to do an automatic patcher to this, because it would be not a smooth solution, but I wonder if anyone know a better way to this problem: to compile AMD_IL code and execute it with the latest OpenCL.

 

I think both CAL and OCL uses the same compiler, but maybe OCL is passing it a special parameter and thus it generates slightly different code.

I hope there is a new enum for that replaces CAL_LANGUAGE_IL. ( calclCompile(obj, CAL_LANGUAGE_IL, PAnsiChar(AIL), ATarget) )

 

(Attaching dumped elf images, sources for the two methods (the thing in !ocl_ocl\ works and the one in !ocl_il fails). You can see, both .il files are the same, but the .isa files has one byte difference)

 

Anyone has some experience with this?

Or is there a way to access the AMD_IL compiler from within the OpenCL runtime itself (that would be a jackpot ).

 

If AMD guys read it: It would be awesome if you could officially enable amd_il and binary_isa as an option to create OpenCL kernels out of.

 

Thanks for replies in advance!

Attachments

Outcomes