Return data from kernel function

[Ajmeer_Khan]

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

[GouthamK_Intel]

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]

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]

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 · Accepted solution]

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

[GouthamK_Intel]

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