GPU HD4600 opencl kernel problem

[SDyul]

Hi, i am compiling offline spir kernel.

When i use it on HD4600 GPU i get the following when I invoke clBuildProgram

error: IGILTargetLowering::Call(): unhandled function call!

Call made to: _Z13get_global_idj()
0x7c53480: i64 = GlobalAddress<i64 (i32)* @_Z13get_global_idj> 0 [ORD=1]
error: midlevel compiler failed build.

The same kernel works fine on amd gpu and on intel cpu. Also works fine if the kernel is compiled as spir64

kernel void kernel_main(const global read_only uint8_t* rgb_t, global write_only uint8_t* grayscale, const image_kernel_info src, const image_kernel_info dst)
{
    uint32_t x = get_global_id(0);
    uint32_t y = get_global_id(1);

    if (is_in_interior(&src, x, y))
    {
        const global rgb* rgb_ = (const global rgb*) sample_2d_clamp(rgb_t, &src, x, y, sizeof( rgb ) );
        float  r = rgb_->r / 255.0f;
        float  g = rgb_->g / 255.0f;
        float  b = rgb_->b / 255.0f;

        float  gray = 0.2989f * (r  * r) + 0.5870f * (g *  g) + 0.1140f * (b  *  b);
        uint8_t  gray_quantized = (uint8_t ) (sqrt(gray) * 255.0f);

        write_2d_uint8( grayscale, &dst, x, y, gray_quantized);
    }
}

[Robert_I_Intel]

Hi Stefan,

Could you please provide the whole file, not just the kernel? You have a number of things used within the kernel that are not defined.

[SDyul]

Hi,

I have uploaded the project here:

https://github.com/kingofthebongo2008/dare12_opencl/tree/master/src/opencl

You can check it out.

Also the kernel works fine if i do not use get_global_id

Seems to me it just cannot find the symbols for 32bit spirs. Seems to me that spir32 bit code, generates function signatures for spir64.

The problem most probably is in the spir clang version.

[Robert_I_Intel]

Hi Stefan,

I think I realized what is going on: you are probably trying to feed SPIR32 binary to run it on the Intel GPU. Please read my article on SPIR https://software.intel.com/en-us/articles/using-spir-for-fun-and-profit-with-intel-opencl-code-builder - it will explain what you need to do. Basically, if you have a 32-bit program, you will need SPIR32 binary to run on CPU OpenCL and SPIR64 binary to run on the GPU OpenCL. Feeding SPIR32 binary to GPU will produce the error you describe.

[SDyul]

I thought these things were supported 32 host / 32 device.

thanks for your help. I moved to 64bit everything and now seems to be fine on intel.

[Oleg_G_]

Hi Robert,

Got the exact error on HD4600. The same code works fine on HD530. I'm using SPIR32 and -device=gpu. SPIR64 doesn't work with HD530 (clBuildProgram failed with error code CL_INVALID_BINARY).

[Oleg_G_]

Hi Robert,

Robert I. (Intel) wrote:

I think I realized what is going on: you are probably trying to feed SPIR32 binary to run it on the Intel GPU. Please read my article on SPIR https://software.intel.com/en-us/articles/using-spir-for-fun-and-profit-with-intel-opencl-code-builder - it will explain what you need to do. Basically, if you have a 32-bit program, you will need SPIR32 binary to run on CPU OpenCL and SPIR64 binary to run on the GPU OpenCL. Feeding SPIR32 binary to GPU will produce the error you describe.

It seems to me that rules have changed. Or maybe this is a bug? It is possible to run win32 application on HD 530 only if opencl binary was made with SPIR32 flag. In case of SPIR64 clBuildProgram fails with error code CL_INVALID_BINARY.

At the same time on HD 4600 win32 application works only with SPIR64 binary. SPIR32 produce error as described in first message of this thread.

It is very inconvenient. How should i select SPIR binary according to GPU adapter? Besides, AMD can utilize both SPIR32 and SPIR64 no matter for what platform application was written for. Why Intel can't do the same?

[Oleg_G_]

Hello Intel,

Is that really so:

Intel monitors this forum Monday – Friday, 09:00 – 17:00 Pacific Time (GMT –7:00). Depending on the amount of research we need to do to track down the answer, it may take a day or two for us to respond.

Desperately need the answer to my question.

[Jeffrey_M_Intel1]

Sorry for the delay.  It can be easier to see when answers are needed for new threads than continuations of old ones.  Definitely understand that inconsistent behavior would be frustrating.  I hope to have more details tomorrow.

[Jeffrey_M_Intel1]

Here is at least part of an answer:

For 5th Generation Core (Broadwell) and forward, SPIR should work as expected in 32 and 64 bit mode.  Use 32 bit SPIR with 32 bit executables and 64 bit SPIR with 64 bit executables.  Going forward this expected pattern should continue.

However, for earlier platforms before full SVM support there are some exceptions:

• For 4th Generation Core (Haswell): only 64 bit SPIR is supported
• For Atom processors so far: only 32 bit SPIR is supported

CL_DEVICE_ADDRESS_BITS should be a reliable indicator of what is supported.  This will report 32 or 64 based on your system.  This should include the behavior described above.  For example, on Haswell processors CL_DEVICE_ADDRESS_BITS should always return 64, on Broadwell and newer it will reflect how the executable was compiled.

If you generate 32 and 64 bit SPIR, applications can check CL_DEVICE_ADDRESS_BITS and load the version indicated.

Are you seeing any behavior that does not fit this pattern?

[Oleg_G_]

Jeffrey M. (Intel) wrote:

Are you seeing any behavior that does not fit this pattern?

It works fine for me. But I tried it only on 530 and 4600. I hope that proposed method is universal. Thanks for the help!