Intel® oneAPI Data Parallel C++
Support for Intel® oneAPI DPC++ Compiler, Intel® oneAPI DPC++ Library, Intel® DPC++ Compatibility Tool, and GDB*
Announcements
Welcome to the Intel Community. If you get an answer you like, please mark it as an Accepted Solution to help others. Thank you!

Return data from kernel function

Ajmeer_Khan
Employee
659 Views

Hi @GouthamK_Intel ,

I want to access a global variable(data) from kernel function(add)  and return the data from add function to parallel_for kernel function. I attached the code below. Please help me  to solve this issue.

 

Regards,

Ajmeer Khan M

0 Kudos
1 Solution
GouthamK_Intel
Moderator
576 Views

Hi Ajmeer Khan,

As mentioned earlier, the variables which are declared outside the kernel are READ-ONLY variables and you cannot assign values inside the kernel without creating buffer and accessor to it. Hence you are getting the error when you try to assign the value to the variable auto CRC;

Instead, create a new local variable and assign the returned value to the local variable.

Change the below line in your code 

From

 

 CRC=ab((int *)&xxx[0]);

 

to

 

auto CRC1=ab((int *)&xxx[0]);

 

 

And, if you wanted to use the returned value outside the kernel, create buffer and accessor to the variable CRC and assign the returned value to that variable inside the kernel.

Please refer the below working code, if you are looking answer for the above scenario. (Refer the lines which are appended with comment "Edited line")

 

#include<CL/sycl.hpp>
#include<iostream>
using namespace std;
using namespace sycl;
int ab(int *);
int main()
{
        int crc=10;
        int x[5]={1,2,3,4,5};
        int len=5;
        int CRC = crc;   //Edited line
        cpu_selector device_selector;
        queue q(device_selector);
        {
                buffer<int,1>xx (x,range<1> (len));
                buffer<int,1>crc_buff (&CRC, range<1>{1}); //Edited line
                q.submit([&](handler &h)
                {
                 auto xxx=xx.get_access<access::mode::discard_write>(h);
                 auto crc_acc=crc_buff.get_access<access::mode::write>(h); //Edited line
                     h.single_task([=]()
                        {
                          crc_acc[0] = ab((int *)&xxx[0]); //Edited line
                         });
                 });
          }
        cout<<"crc"<<CRC; //Edited line
}
int ab(int * a)
{
        return a[2];
}

 

Please let us know if this helps.

Have a Good day!

 

Thanks & Regards

Goutham

 

View solution in original post

5 Replies
GouthamK_Intel
Moderator
637 Views

Hi Ajmeer Khan,

We have gone through the source code which you have provided. You are trying to modify the global variable (int data) inside the kernel function(add). Hence you are facing the errors.

You can only use variables (Local/Global) directly as read-only variables inside the kernel. But If you wanted to modify/change the value of those variables, you need to explicitly create buffers/accessors for those variables or use the USM model.

Also, for global variables, you need to declare the variable as const.

Please refer to the below-modified working code of yours.

#include<iostream>
#include<CL/sycl.hpp>

using namespace std;
using namespace sycl;

const int mydata=100;

int add(int *p, int x);

int main()
{
    int a[10],b[10],c[10];
    for(int i=0;i<10;i++)
    {
        a[i]=i+1;
        b[i]=2;

    }

    cpu_selector device_selector;
    queue q(device_selector);
    {
        buffer<int ,1> a_buff(a,range<1> (10));
       // buffer<int ,1> b_buff(b,range<1> (10));
        //buffer<int,1> c_buff(c,range<1> (10));
        cout<<q.get_device().get_info<info::device::name>().c_str() << "\n";
        q.submit([&](handler &h)
        {
            cl::sycl::stream kernelout(512*1024, 512, h);
            auto aa=a_buff.get_access<access::mode::write>(h);
           // auto bb=b_buff.get_access<access::mode::write>(h);
            //auto cc=c_buff.get_access<access::mode::write>(h);
            h.parallel_for(range<1>{10},[=](id<1> item)
            {
                int result = add((int *)&aa[0],3);
                aa[6]=result;
            });
      });
    }
    for(int i=0;i<10;i++)
    {
        cout<<a[i]<<"\t";
    }

}


int add(int *p,int x)
{
    p[x]=80;
    int l_data = x * mydata;
    p[x+1]=l_data;
    return l_data;
}

Let us know if this helps and whether your issue is resolved.

Have a Good day!

 

Thanks & Regards

Goutham

 

GouthamK_Intel
Moderator
613 Views

Hi Ajmeer Khan,

Could you please let us know if your issue is resolved?

If your issue still persists, let us know the issue you are facing.


Thanks & Regards

Goutham


Ajmeer_Khan
Employee
598 Views

Hi Goutham,

I am still facing an issue with return data.. How do I assign the return data(a[2])  into accessor(CRC).

Regards,

Ajmeer khan M

GouthamK_Intel
Moderator
577 Views

Hi Ajmeer Khan,

As mentioned earlier, the variables which are declared outside the kernel are READ-ONLY variables and you cannot assign values inside the kernel without creating buffer and accessor to it. Hence you are getting the error when you try to assign the value to the variable auto CRC;

Instead, create a new local variable and assign the returned value to the local variable.

Change the below line in your code 

From

 

 CRC=ab((int *)&xxx[0]);

 

to

 

auto CRC1=ab((int *)&xxx[0]);

 

 

And, if you wanted to use the returned value outside the kernel, create buffer and accessor to the variable CRC and assign the returned value to that variable inside the kernel.

Please refer the below working code, if you are looking answer for the above scenario. (Refer the lines which are appended with comment "Edited line")

 

#include<CL/sycl.hpp>
#include<iostream>
using namespace std;
using namespace sycl;
int ab(int *);
int main()
{
        int crc=10;
        int x[5]={1,2,3,4,5};
        int len=5;
        int CRC = crc;   //Edited line
        cpu_selector device_selector;
        queue q(device_selector);
        {
                buffer<int,1>xx (x,range<1> (len));
                buffer<int,1>crc_buff (&CRC, range<1>{1}); //Edited line
                q.submit([&](handler &h)
                {
                 auto xxx=xx.get_access<access::mode::discard_write>(h);
                 auto crc_acc=crc_buff.get_access<access::mode::write>(h); //Edited line
                     h.single_task([=]()
                        {
                          crc_acc[0] = ab((int *)&xxx[0]); //Edited line
                         });
                 });
          }
        cout<<"crc"<<CRC; //Edited line
}
int ab(int * a)
{
        return a[2];
}

 

Please let us know if this helps.

Have a Good day!

 

Thanks & Regards

Goutham

 

View solution in original post

GouthamK_Intel
Moderator
563 Views

Hi, 

Thanks for the confirmation.

Glad to know that your issue is resolved!

As this issue has been resolved, we will no longer respond to this thread. 

If you require any additional assistance from Intel, please start a new thread. 

Any further interaction in this thread will be considered community only. 

Have a Good day!


Regards

Goutham


Reply