topic Hi Kamil, in Intel® oneAPI DPC++/C++ Compiler

Error -50 (CL_INVALID_ARG_VALUE) for Intel iGPU

H__Kamil — Wed, 08 Apr 2020 21:08:45 GMT

Hello,
I have a problem with execution of the following code:

#include <CL/sycl.hpp>
#include <iostream>
#include <array>
#include <cstdio>
using namespace std;
using namespace cl::sycl;

#define SIZE 20

struct DeviceData
{
    queue q;
    
    int* A;
    int* d_A;
    
    void init()
    {
        q = queue(gpu_selector{});
        
        A = (int* )malloc(sizeof(int)*SIZE);
        for(int i=0; i<SIZE; ++i)
            A = i+1;

        d_A = (int *)malloc_device(sizeof(int)*SIZE, q.get_device(), q.get_context());
    }
     
    void test()
    {
        cout << "Before" << "\n";
        for(int i=0; i<SIZE; ++i)
        {
            cout << A << " ";
        }
        cout << "\n";
   
        q.submit([&](handler &h){
            h.memcpy(d_A, A, sizeof(int)*SIZE);
        });
        q.wait();

        q.submit([&](handler &h){
            h.parallel_for(range<1>{SIZE}, [=](id<1> i){
                d_A += 10;
            });
        });
        q.wait();
       
        q.submit([&](handler &h){
            h.memcpy(A, d_A, sizeof(int)*SIZE);
        });
        q.wait();
        
        cout << "After" << "\n";
        for(int i=0; i<SIZE; ++i)
            cout << A << " ";
        cout << "\n";
    }
    void free()
    {
       free(A);
       free(d_A, q.get_context());
    }
};

int main()
{
    DeviceData dev;
    dev.init();
    dev.test(); 
    dev.free();
    return 0;
}

The above code works well for CPU (when using cpu_selector{}), while it fails for GPU. I obtain the error: OpenCL API failed. OpenCL API returns: -50 (CL_INVALID_ARG_VALUE) -50 (CL_INVALID_ARG_VALUE). The problem occurs during execution of the kernel within test() function. This code is used to present the problem. I am working with more complex application which is implemented in the presented manner (using structures). The code is executed on Intel DevCloud.

I will be grateful for any advice.
Thanks :)

Hi Kamil,

GouthamK_Intel — Thu, 09 Apr 2020 11:11:00 GMT

Hi Kamil,

Thanks for reaching out to us!

Could you please attach the error logs or screenshots of the output for the above code after running on devcloud and also mention the node name on which you are running. So that we would be able to investigate more on your issue.

Please provide the logs for both the cases after running your code with cpu_selector and gpu_selector.

Thanks

Goutham.

Output for cpu_seletor:

H__Kamil — Thu, 09 Apr 2020 13:43:00 GMT

Output for cpu_seletor:

./main.exe
Before
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20
After
11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30

Output for gpu_selector:

./main.exe
Before
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20
terminate called after throwing an instance of 'cl::sycl::runtime_error'
  what():  OpenCL API failed. OpenCL API returns: -50 (CL_INVALID_ARG_VALUE) -50 (CL_INVALID_ARG_VALUE)
Aborted

Node name: s001-n160

Thanks for help!

Hi Kamil,

GouthamK_Intel — Mon, 13 Apr 2020 11:17:17 GMT

Hi Kamil,

We are able to reproduce the error which you are facing.

Please find the below workaround for your code.

#include <CL/sycl.hpp>
#include <iostream>
#include <array>
#include <cstdio>
using namespace std;
using namespace cl::sycl;

#define SIZE 20

struct DeviceData
{
    queue q;
    
    int* A;
    int* d_A;

    void init(int* A,int* d_A,queue q)
    {
        q = queue(gpu_selector{});
        
//        A = (int* )malloc(sizeof(int)*SIZE);
        for(int i=0; i<SIZE; ++i)
            A = i+1;

  //      d_A = (int *)malloc_device(sizeof(int)*SIZE, q.get_device(), q.get_context());
    }
     
    void test(int* A,int* d_A,queue q)
    {
        cout << "Before" << "\n";
        for(int i=0; i<SIZE; ++i)
        {
            cout << A << " ";
        }
        cout << "\n";
   
        q.submit([&](handler &h){
            h.memcpy(d_A, A, sizeof(int)*SIZE);
        });
        q.wait();

        q.submit([&](handler &h){
            h.parallel_for(range<1>{SIZE}, [=](id<1> i){
                d_A += 10;
            });
        });
        q.wait();
       
        q.submit([&](handler &h){
            h.memcpy(A, d_A, sizeof(int)*SIZE);
        });
        q.wait();
        
        cout << "After" << "\n";
        for(int i=0; i<SIZE; ++i)
            cout << A << " ";
        cout << "\n";
    }
    void free(int* A,int* d_A,queue q)
    {
       std::free(A);
       cl::sycl::free(d_A, q.get_context());
    }
};


int main()
{
    DeviceData dev;
    dev.d_A=(int *)malloc_device(sizeof(int)*SIZE, dev.q.get_device(), dev.q.get_context());
    
    dev.A=(int* )malloc(sizeof(int)*SIZE);
    dev.init(dev.A,dev.d_A,dev.q);
    dev.test(dev.A,dev.d_A,dev.q); 
    dev.free(dev.A,dev.d_A,dev.q);
    return 0;
}

Please let us know if this resolves your issue.

Regards

Goutham

Hi Kamil,

GouthamK_Intel — Thu, 16 Apr 2020 05:30:13 GMT

Hi Kamil,

Could you please let us know if your issue is resolved or if you are facing any issues with the code?

Thanks

Goutham

Hi,

H__Kamil — Thu, 16 Apr 2020 17:46:59 GMT

Hi,

i am apologize for delay - spring holiday. Yes, the code provided by you works fine.

But, I have a question. Why my code does not work properly? I don't know if I am right, but is seems that the queue should be created within the same instruction block next to kernel and buffers.

Best regards,
Kamil

Hi Kamil,

GouthamK_Intel — Wed, 22 Apr 2020 06:28:52 GMT

Hi Kamil,

We are investigating more about your issue.

We are working with the concerned team and we will get back to you with more information.

Regards

Goutham

Hi,

H__Kamil — Wed, 22 Apr 2020 08:00:51 GMT

Hi,

i am waiting for more information :) Thanks you for helping me solve the problem :)

Best regards

Hi Kamil,

Anoop_M_Intel — Wed, 22 Apr 2020 12:02:00 GMT

Hi Kamil,

In your program, d_A is a device pointer. When accessing the class member data within the class member functions, the "this" pointer is passed implicitly and that is used to access the object members. So in this case, since d_A is a class member, any reference to d_A inside the class member function will be resolved as this->d_A. The caveat is, "this" pointer points to the memory where the class object is located and this class object is created on CPU memory. Hence when we try to use d_A inside a DPC++ kernel which is supposed to run inside a GPU, the expression "this->d_A" is invalid since "this" is invalid pointer from GPU side. The same code works on the CPU side since "this" pointer is valid pointer when the code executes on CPU. Below is a way to work around this issue by capturing the "this->d_A - device pointer" in a local variable in the kernel function (also I have modified the code to introduce newer syntax for USM which decreases the code verbosity):

#include <CL/sycl.hpp>
#include <iostream>
#include <array>
#include <cstdio>
using namespace std;
using namespace cl::sycl;

#define SIZE 20

struct DeviceData
{
    queue q;

    int* A;
    int* d_A;

    void init()
    {
        q = queue(gpu_selector{});

        A = (int* )malloc(sizeof(int)*SIZE);
        for(int i=0; i<SIZE; ++i)
            A = i+1;

        d_A = (int *)malloc_device(sizeof(int)*SIZE, q.get_device(), q.get_context());
    }

    void test()
    {
        cout << "Before" << "\n";
        for(int i=0; i<SIZE; ++i)
        {
            cout << A << " ";
        }
        cout << "\n";

        q.memcpy(d_A, A, sizeof(int)*SIZE);
        q.parallel_for(range<1>{SIZE}, [=,d_A_local=this->d_A](id<1> i){
                d_A_local += 10;
        });
        q.memcpy(A, d_A, sizeof(int)*SIZE).wait();

        cout << "After" << "\n";
        for(int i=0; i<SIZE; ++i)
            cout << A << " ";
        cout << "\n";
    }
    void free()
    {
            std::free(A);
            sycl::free(d_A, q.get_context());
    }
};

int main()
{
    DeviceData dev;
    dev.init();
    dev.test();
    dev.free();
    return 0;
}

Hi Anoop,

GouthamK_Intel — Thu, 23 Apr 2020 07:09:27 GMT

Hi Anoop,

Thanks for providing a detailed explanation.

Hi Kamil,

Please let us know if the information provided is helpful.

Confirm if your issue is resolved and let us know whether we can close this thread.

Thanks

Goutham

Hi Kamil,

GouthamK_Intel — Tue, 28 Apr 2020 06:03:45 GMT

Hi Kamil,

Please confirm if the explanation provided helped.

Let us know if we can close this thread.

Regards

Goutham

Hi,

H__Kamil — Tue, 28 Apr 2020 08:14:09 GMT

Hi,

yes, this advice resolve my problem completely.

I have another question, but I don't want to create new topic. I have a problem with math functions like cos(), sin() etc. on the iGPU side. When I am trying use them within the kernel I obtain the error: undefined reference to `cos()'.

Best regards,
Kamil

Hi Kamil,

GouthamK_Intel — Tue, 28 Apr 2020 10:05:09 GMT

Hi Kamil,

We are glad to know that the solution provided resolved your issue.

Sure, we will help you with your other issue. Please provide your code if you can. So that it will help us to investigate.

Regards

Goutham

I modified code presented in

H__Kamil — Tue, 28 Apr 2020 10:49:43 GMT

I modified code presented in this topic to show the problem:

#include <CL/sycl.hpp>
#include <iostream>
#include <array>
#include <cstdio>
#include <cmath>
using std::cout;
using std::endl;
using namespace cl::sycl;

#define SIZE 20

struct DeviceData
{
    queue q;

    double* A;
    double* d_A;

    void init()
    {
        q = queue(gpu_selector{});

        A = (double* )malloc(sizeof(double)*SIZE);
        for(int i=0; i<SIZE; ++i)
            A = i+1;

        d_A = (double *)malloc_device(sizeof(double)*SIZE, q.get_device(), q.get_context());
    }

    void test()
    {
        cout << "Before" << "\n";
        for(int i=0; i<SIZE; ++i)
        {
            cout << A << " ";
        }
        cout << "\n";

        q.memcpy(d_A, A, sizeof(double)*SIZE);
        q.parallel_for(range<1>{SIZE}, [=,d_A_local=this->d_A](id<1> i){
                d_A_local = cos(i+1);
        });
        q.memcpy(A, d_A, sizeof(double)*SIZE).wait();

        cout << "After" << "\n";
        for(int i=0; i<SIZE; ++i)
            cout << A << " ";
        cout << "\n";
    }
    void free()
    {
            std::free(A);
            sycl::free(d_A, q.get_context());
    }
};

int main()
{
    DeviceData dev;
    dev.init();
    dev.test();
    dev.free();
    return 0;
}

Output:

./main.exe
Before
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20
terminate called after throwing an instance of 'cl::sycl::compile_program_error'
  what():  The program was built for 1 devices
Build program log for 'Intel(R) Gen9 HD Graphics NEO':

error: undefined reference to `cos()'

error: backend compiler failed build.
 0 (CL_SUCCESS)
Makefile:43: recipe for target 'run' failed
make: *** [run] Aborted

Hi Kamil,

GouthamK_Intel — Wed, 29 Apr 2020 12:06:32 GMT

Hi Kamil,

Please find the attached code which will resolve the issue. As per the SYCL standard, the built-in functions (sin(),cos()..etc) can take as input float or optionally double.

#include <CL/sycl.hpp>
#include <iostream>
#include <array>
#include <cstdio>
#include <cmath>
using std::cout;
using std::endl;
using namespace cl::sycl;

#define SIZE 20

struct DeviceData
{
    queue q;

    double* A;
    double* d_A;

    void init()
    {
        q = queue(gpu_selector{});

        A = (double* )malloc(sizeof(double)*SIZE);
        for(int i=0; i<SIZE; ++i)
            A = i+1;

        d_A = (double *)malloc_device(sizeof(double)*SIZE, q.get_device(), q.get_context());
    }

    void test()
    {
        cout << "Before" << "\n";
        for(int i=0; i<SIZE; ++i)
        {
            cout << A << " ";
        }
        cout << "\n";

        q.memcpy(d_A, A, sizeof(double)*SIZE);
        q.parallel_for(range<1>{SIZE}, [=,d_A_local=this->d_A](id<1> i){
                d_A_local = cl::sycl::cos((float)i+1); // cos() takes input as float datatype
        });
        q.memcpy(A, d_A, sizeof(double)*SIZE).wait();

        cout << "After" << "\n";
        for(int i=0; i<SIZE; ++i)
            cout << A << " ";
        cout << "\n";
    }
    void free()
    {
            std::free(A);
            sycl::free(d_A, q.get_context());
    }
};

int main()
{
    DeviceData dev;
    dev.init();
    dev.test();
    dev.free();
    return 0;
}

Please confirm if this resolves the issue.

Regards

Goutham

Hi,

H__Kamil — Fri, 01 May 2020 14:47:41 GMT

Hi,

yes, it works fine. Thanks for help :)

Hi Kamil,

GouthamK_Intel — Mon, 04 May 2020 05:55:18 GMT

Hi Kamil,

Glad to hear that your issue got resolved.

We are closing this thread.

Feel free to raise a new thread in case of any further support.

Regards

Goutham