"One of" or "choice" kernel

[AdamVerilab]

Hello,

I was wondering how to model a choice in DPC++.

In other words, a number of kernels are scheduled but any one of them actually being chosen for execution means that the others don't need to run.

The various kernels in question would share the same input and output buffers.

Thanks,

Adam.

[RahulV_intel]

Hi Adam,

A simple approach would be to use conditional directives(#ifdef, #elif) and define your kernels under these conditional statements.

PFA attached code below.

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

using namespace sycl;

#define size 1024

int main() {
	std::vector<int> A(size, size), B(size, size), C(size, 0);

 { //SYCL scope begins
	queue q;
	range<1> R(size);
	buffer<int,1> buffA(A.data(), R);
	buffer<int,1> buffB(B.data(), R);
	buffer<int,1> buffC(C.data(), R);
	#ifdef kernel1
		std::cout<<"Running on: "<<q.get_device().get_info<sycl::info::device::name>()<<"\n";
		q.submit([&](handler &cgh) {
			auto acc_buffA = buffA.get_access<access::mode::read>(cgh);
			auto acc_buffB = buffB.get_access<access::mode::read>(cgh);
			auto acc_buffC = buffC.get_access<access::mode::write>(cgh);
			cgh.parallel_for(R, [=](id<1> i) {
				acc_buffC[i] = acc_buffA[i] + acc_buffB[i];
			});
		});
	#elif kernel2
		std::cout<<"Running on: "<<q.get_device().get_info<sycl::info::device::name>()<<"\n";
		q.submit([&](handler &cgh) {
			auto acc_buffA = buffA.get_access<access::mode::read>(cgh);
			auto acc_buffB = buffB.get_access<access::mode::read>(cgh);
			auto acc_buffC = buffC.get_access<access::mode::write>(cgh);
			cgh.parallel_for(R, [=](id<1> i) {
				acc_buffC[i] = acc_buffA[i] + acc_buffB[i];
			});
		});
	#endif
	} //SYCL scope ends

	 std::vector<int> vecValidate(size, 2*size);
	(C==vecValidate) ? std::cout << "Success\n" : std::cout<<"Failure\n";
	return 0;
}

Here, I have defined two similar kernels (vector add) as kernel1 and kernel2. During compilation, you may either specify kernel1 or kernel2 as an argument to compute any one kernel.

#To compile and run kernel1
dpcpp vecadd.cpp -Dkernel1 -o k1add && ./k1add

#To compile and run kernel2
dpcpp vecadd.cpp -Dkernel2 -o k2add && ./k2add

Let us know if this is the kind of arrangement that you are looking for. Else, please elaborate more on your exact requirements.

Thanks,

Rahul

[AdamVerilab]

Hi Rahul, thanks for your reply. 

A compile time approach is not what I am looking for. I was thinking of genuine run time non determinacy. The first kernel to run would somehow turn off the other kernels in the set.

In a way, this is question about how the kernel scheduler works. For example, if the first thing any of these kernels did is to set a flag which contributed to an any_of vote, would this be seen soon enough to turn off the other kernels ? Or might one of the other parallel kernels already running ?

[RahulV_intel]

Hi Adam,

A simple runtime approach would be to take the input kernel (as input at runtime) that you wish to run and use an if/else condition accordingly.

To read more about kernel scheduling and other DPC++ concepts, you may refer to the DPC++ book given in the link below:

https://www.apress.com/us/book/9781484255735'

Let us know if it helps.

Thanks,

Rahul

[RahulV_intel]

Hi Adam,

Could you please confirm if your issue is resolved? Let us know if you have any queries.

Thanks,

Rahul

[RahulV_intel]

Hi,

I have not heard back from you, so I will go ahead and close this thread from my end. Intel will no longer monitor this thread. Feel free to post a new query if you require further assistance from Intel.