- Mark as New
- Bookmark
- Subscribe
- Mute
- Subscribe to RSS Feed
- Permalink
- Email to a Friend
- Report Inappropriate Content

I have 2 vectors (float) with size of 1024*1024*8. I want to do vector addition. My first kernel vec_add_1() has Gx=1024*1024*8 and Lx=0

__kernel void vec_add_1(__global const float* in1, __global const float* in2, __global float* out) {

int i=get_global_id(0);

out*=in1 +in2;*

}

Kernel vec_add_1() takes about 10msec.

To reduce schedule time, I created second kernel vec_add_2(). vec_add_1() has Gx=1024*1024*8 /4, Lx=0.

__kernel void vec_add_2(__global const float* in1, __global const float* in2, __global float* out) {

int i=get_global_id(0);

int j=(i<<2);

out

out[j+1]=in1[j+1]+in2[j+1];

out[j+2]=in1[j+2]+in2[j+2];

out[j+3]=in1[j+3]+in2[j+3];

}

However, I got 2 quite different results

- Running vec_add_2() in code builder session, vec_add_2() takes ~13msec, which is slower than vec_add_1()

- Running vec_add_2() with host code together, vec_add_2() takes ~7msec, which is faster than vec_add_1()

So my questions are

- why running vec_add_2() with and without code builder session give quite different results?

- is vec_add_2() an optimized version than vec_add_1()?

thanks,

Jeffrey

Link Copied

- Mark as New
- Bookmark
- Subscribe
- Mute
- Subscribe to RSS Feed
- Permalink
- Email to a Friend
- Report Inappropriate Content

Hi Jeffrey,

I think it would be a good idea to see my presentations on Optimizing Simple OpenCL Kernels here https://software.intel.com/en-us/articles/optimizing-simple-opencl-kernels - there is also code there that you can download and play with. Basically, what you are doing is not very optimal. You are better off switching to float4 datatype.

- Mark as New
- Bookmark
- Subscribe
- Mute
- Subscribe to RSS Feed
- Permalink
- Email to a Friend
- Report Inappropriate Content

Hi Robert,

I changed my kernel to use float4 as below

__kernel void vec_add_3(__global const float4* in1, __global const float4* in2, __global float4* out)

{

int i=get_global_id(0);

out*=in1 +in2;*

}

Gx is correspondingly changed to be 1/4 of original vector size (original vector size is 1024*1024*8. Gx for vec_add_3() is 1024*1024*8 / 4). Then I did "run analysis" in code-builder session. Surprisingly, I execution time is still ~10ms, almost no improvement.

- Mark as New
- Bookmark
- Subscribe
- Mute
- Subscribe to RSS Feed
- Permalink
- Email to a Friend
- Report Inappropriate Content

Hi Jeffrey,

Couple of more things two try:

1. try to combine what you are doing in vec_add_2 with float4

2. try using float8 or even float16

3. try a combination of vec_add_2 and float4 and/or float8.

The basic problem is that for the simple kernels you need to pack much more compute onto a hardware thread.

- Subscribe to RSS Feed
- Mark Topic as New
- Mark Topic as Read
- Float this Topic for Current User
- Bookmark
- Subscribe
- Printer Friendly Page