UPGRADE YOUR BROWSER

We have detected your current browser version is not the latest one. Xilinx.com uses the latest web technologies to bring you the best online experience possible. Please upgrade to a Xilinx.com supported browser:Chrome, Firefox, Internet Explorer 11, Safari. Thank you!

cancel
Showing results for 
Search instead for 
Did you mean: 
Highlighted
Observer race
Observer
764 Views
Registered: ‎05-18-2017

Accessing uninitialized memory -- works on CPU emulation.

Hi,

 

 

The following code is an optimized opencl implementation on GPU. I am trying to run the same code on SDAccel but the HW emulation gives lots of warning that I am trying to access uninitialized memory. The output is not correct, but the functionality is right on CPU emulation.

 

// TODO: Add OpenCL kernel code here.
__kernel void gaus(__global uchar16* input, __global uchar16* output)
{
//Opencl constructs to get program parameters
uint x = get_global_id(0);
uint y = get_global_id(1);
uint width = get_global_size(0);
uint height = get_global_size(1);
//calculating source and destination index
uint dst_index = 16 * y * width + x;
/* src size is increased by 16 and border added */
uint src_index = 16 * y * (width + 2) + x + 1;

//uint src_index = y * width + x;
//perform calculation
float i0; float16 i1; float i2;
float i3; float16 i4; float i5;
float i6; float16 i7; float i8;
float16 Gx;
//pixel data
i0 = convert_float(((__global uchar*)(input+src_index))[-1]);
//printf("%f",i0);
//i0 = convert_float((__global uchar*)(input+src_index)[-1]); 
i1 = convert_float16(input[src_index]); 
i2 = convert_float(((__global uchar*)(input+src_index))[16]);
src_index += width + 2;
i3 = convert_float(((__global uchar*)(input+src_index))[-1]);
i4 = convert_float16(input[src_index]); 
i5 = convert_float(((__global uchar*)(input+src_index))[16]);

for(int i = 0; i < 16; i++)
{
src_index += width + 2;
//collect new data
i6 = convert_float(((__global uchar*)(input+src_index))[-1]);
i7 = convert_float16(input[src_index]); 
i8 = convert_float(((__global uchar*)(input+src_index))[16]);
Gx = ((float16)(i0, i1.s0123, i1.s456789ab, i1.scde)
+ 2.0f * i1
+ (float16)(i1.s123, i1.s4567, i1.s89abcdef, i2)
+ 2.0f * (float16)(i3, i4.s0123, i4.s456789ab, i4.scde)
+ 4.0f * i4
+ 2.0f * (float16)(i4.s123, i4.s4567, i4.s89abcdef, i5)
+ (float16)(i6, i7.s0123, i7.s456789ab, i7.scde)
+ 2.0f * i7
+ (float16)(i7.s123, i7.s4567, i7.s89abcdef, i8)) / 16.0f;
output[dst_index] = convert_uchar16_sat(Gx);
//reuse old data
i0 = i3; i3 = i6;
i1 = i4; i4 = i7;
i2 = i5; i5 = i8;
dst_index += width;
}
 
}

 

Tags (2)
0 Kudos
1 Reply
Observer race
Observer
759 Views
Registered: ‎05-18-2017

Re: Accessing uninitialized memory -- works on CPU emulation.

host side image padding

 

copyMakeBorder(input_image, image_pad ,1 ,1 ,1 ,1 ,BORDER_REPLICATE, Scalar(0));
copyMakeBorder(image_pad, image_pad ,0 ,0 ,15 ,15 ,BORDER_CONSTANT, Scalar(0));

0 Kudos