# Concurrents kernels in OpenCL

**URL:** <https://community.khronos.org/t/concurrents-kernels-in-opencl/7604>\
**Category:** OpenCL\
**Created:** [November 16, 2018, 7:32am UTC](https://community.khronos.org/t/concurrents-kernels-in-opencl/7604 "2018-11-16T07:32:07Z")\
**Posts on this page:** 5\
**Page:** 1

<div class="post-metadata">

**Author:** ![hanarce](https://community.khronos.org/letter_avatar_proxy/v4/letter/h/53a042/32.png) [@hanarce](https://community.khronos.org/u/hanarce)\
**Post date:** [November 16, 2018, 7:32am UTC](https://community.khronos.org/t/concurrents-kernels-in-opencl/7604/1 "2018-11-16T07:32:07Z")

</div>

Hello, I would like know how can I execute two or more different kernels in parallel and at the same time? Obviously in the same GPU using OpenCL.  
My main idea is to use two different kernels (kernel A and Kernel B) but they need to use the same memory (I do not want to duplicate the memory by using one buffer for each in the “a” and “b” pointers). So is there another way I can accomplish the dual execution with an efficient memory technique?  
The codes of the kernels are the following:  
Kernel A:

```auto

_kernel void kernelA(global struct VectorStruct* a, int aLen0, global struct VectorStruct* b, int bLen0, global struct VectorStruct* c, int cLen0)
{
	int i = get_local_id(0);
	c[(i)].x = a[(i)].x + b[(i)].x;
}

```

Kernel B:

```auto

kernel void kernelB(global struct VectorStruct* a, int aLen0, global struct VectorStruct* b, int bLen0, global struct VectorStruct* d, int cLen0)
{
	int i = get_local_id(0);
	d[(i)].y = a[(i)].y + b[(i)].y;
}

```

The definition for the struct VectorStruct is the following:

```auto

struct VectorStruct
{
	int x;
	int y;
};

```

In the host code I have to create four pointers:  
VectorStruct\* a  
VectorStruct\* b  
VectorStruct\* c  
VectorStruct\* d  
The poiner “a” and “b” have the data that I will transfer to GPU. The pointer “c” will storage the results of the kernel A, and the pointer “d” will storage the results of the kernel B.

---

<div class="post-metadata">

**Author:** ![OmegaDoom](https://community.khronos.org/letter_avatar_proxy/v4/letter/o/ed8c4c/32.png) [@OmegaDoom](https://community.khronos.org/u/OmegaDoom)\
**Post date:** [November 16, 2018, 11:22am UTC](https://community.khronos.org/t/concurrents-kernels-in-opencl/7604/2 "2018-11-16T11:22:46Z")

</div>

Hello hanarce.

I don’t see any problem here. clEnqueueNDRangeKernel is asynchronous by nature. You just need to enqueue second kernel right after the first one. Modern gpus can run kernels in parallel but you need to check if it’s your case. Also depends if resources are available.

---

<div class="post-metadata">

**Author:** ![Dithermaster](https://community.khronos.org/letter_avatar_proxy/v4/letter/d/3ab097/32.png) [@Dithermaster](https://community.khronos.org/u/Dithermaster)\
**Post date:** [November 20, 2018, 3:42pm UTC](https://community.khronos.org/t/concurrents-kernels-in-opencl/7604/3 "2018-11-20T15:42:28Z")

</div>

It also depends on if they use the same resources (for example, if kernelA writes to a buffer than kernelB takes, the runtime should not run them in parallel). A given vendor might make you jump through more hoops, such as using separate command queues for kernelA and kernelB.

---

<div class="post-metadata">

**Author:** ![bashbaug](https://community.khronos.org/user_avatar/community.khronos.org/bashbaug/32/3674_2.png) [@bashbaug](https://community.khronos.org/u/bashbaug)\
**Post date:** [November 21, 2018, 11:43am UTC](https://community.khronos.org/t/concurrents-kernels-in-opencl/7604/4 "2018-11-21T11:43:14Z")

</div>

> [@hanarce;43969](#):
>
> I would like know how can I execute two or more different kernels in parallel and at the same time? Obviously in the same GPU using OpenCL.

This is a deceptively complex question :D.

The short answer is that some OpenCL implementations will support this, usually either by using out-of-order queues or multiple in-order queues or some combination of the two.

The Execution Model chapter in the OpenCL spec describes the conditions by which kernel instances may run in parallel. Note, the OpenCL specification currently doesn’t _guarantee_ that multiple kernel instances _will_ run in parallel, but you can submit command such that they _may_ run in parallel.

[https://www.khronos.org/registry/OpenCL/specs/2.2/html/OpenCL\_API.html#\_execution\_model](https://www.khronos.org/registry/OpenCL/specs/2.2/html/OpenCL_API.html#_execution_model)

I’d also recommend reading the Memory Model chapter in the OpenCL spec, which describes how to safely access memory from kernel instances that may run in parallel, particularly if multiple kernel instances will be writing to the same memory.

[https://www.khronos.org/registry/OpenCL/specs/2.2/html/OpenCL\_API.html#\_memory\_model](https://www.khronos.org/registry/OpenCL/specs/2.2/html/OpenCL_API.html#_memory_model)

Good luck!

---

<div class="post-metadata">

**Author:** ![L30nardoSV](https://community.khronos.org/user_avatar/community.khronos.org/l30nardosv/32/4907_2.png) [@L30nardoSV](https://community.khronos.org/u/L30nardoSV)\
**Post date:** [November 22, 2018, 5:08am UTC](https://community.khronos.org/t/concurrents-kernels-in-opencl/7604/5 "2018-11-22T05:08:12Z")

</div>

> [@bashbaug;43988](#):
>
> The short answer is that some OpenCL implementations will support this, usually either by using out-of-order queues or multiple in-order queues or some combination of the two.
> 
> The Execution Model chapter in the OpenCL spec describes the conditions by which kernel instances may run in parallel. Note, the OpenCL specification currently doesn’t _guarantee_ that multiple kernel instances _will_ run in parallel, but you can submit command such that they _may_ run in parallel.

Interesting!

is there a way to identify which devices actually support concurrent kernel execution?

I have been trying using Intel multicore CPUs (that support CL2.0) for a program that communicates via pipes, but it seems that one kernel has to complete execution before others can start:  
[https://software.intel.com/en-us/forums/opencl/topic/559760](https://software.intel.com/en-us/forums/opencl/topic/559760)

Any comments on this?

Leo
