clBuildProgram bug when using image2d_t and driver 260.61?

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello there!

I may have discovered a bug using image2d_t in opencl code. I can get the clBuildProgram() call to crash (AccessViolation) every time when I have a kernel that looks like this:

__kernel void BoxRowsTex( __read_only image2d_t SourceRgbaTex, __global unsigned int* uiDest, sampler_t RowSampler,
unsigned int uiWidth, unsigned int uiHeight, int iRadius, float fScale)
{
}

When I don’t use any image2d_t and/or sampler_t calls in the kernel it causes the clBuildProgram to crash.

And when I do:

__kernel void BoxRowsTex( __read_only image2d_t SourceRgbaTex, __global unsigned int* uiDest, sampler_t RowSampler,
unsigned int uiWidth, unsigned int uiHeight, int iRadius, float fScale)
{
int2 pos = {0 , 0};
f4Sum += convert_float4(read_imageui(SourceRgbaTex, RowSampler, pos));
}

It compiles fine for me. This seems to be because the call “read_imageui(SourceRgbaTex, RowSampler, pos)” does something with image2d_t and sampler_t.

Can someone else please confirm this? You can easily try this by modifying the oclBoxFilter example in the SDK.

I am running Windows XP, Geforce 8400 GS and driver 260.61.

/anlmat

2

I’m having the same problem, but it is not related to image2d_t. I switched to developer driver 260.61 yesterday (on Windows XP).

The code below will crash clBuildProgram() with an access violation. Comment out the kernel’s body (leaving an empty kernel, i.e. {}) and it will still crash, comment out the remaining functions and it will build.

I have two or three different kernels crashing that way.

void rgbToHsv(uint4 rgb, uint* uiH, uint* uiS, uint* uiV)

{  

	float r = rgb.x;

	float g = rgb.y;

	float b = rgb.z;

	float max = max(r, max(g, b));

	float min = min(r, min(g, b));

	float delta = max - min;

	*uiV = (uint) max;

	if (max == 0.0f) // r = g = b = 0

	{

		*uiS = 0;

		*uiH = 0;

		return;

	}

	else 

		*uiS = convert_uint_rte(255.0f * delta / max);

	float fH = 0.0f; // Case: max == min

	

	if (max != min)

	{

		float fInvDelta = 1.0f / delta;

		if (r == max)

			fH = (g - b) * fInvDelta;

		else if (g == max)

			fH = 2.0f + (b - r) * fInvDelta;

		else

			fH = 4.0f + (r - g) * fInvDelta;

	}

	fH *= 60.0f;

		

	if (fH < 0.0f)

		fH += 360.0f;

	

	*uiH = convert_uint_rte(fH);

}

uint hsvPackToUint(uint h, uint s, uint v)

{

	uint i = 0x000000FF & v;

	i |= 0x0000FF00 & (s << 8);

	i |= 0xFFFF0000 & (h << 16);

	return i;

}

__kernel void 

RgbBufToHsvBuf(

	__global __read_only uchar* ucRgbBuf,

	__global __write_only uint* uiHsvBuf,

	uint uiWidth,

	uint uiHeight)

{

	uint globalX = get_global_id(0);

	

	if (globalX >= uiWidth)

	{

		return;

	}

	for (uint y = 0, uiHsvBufIdx = globalX;

		y != uiHeight;

		++y, uiHsvBufIdx += uiWidth)

	{

		uint4 rgba;

		uint offset = uiHsvBufIdx * 3;

		rgba.x = ucRgbBuf[offset];

		rgba.y = ucRgbBuf[offset + 1];

		rgba.z = ucRgbBuf[offset + 2];

		rgba.w = 255;

	

		uint h, s, v;

#ifdef RGB_REVERSED

		rgbToHsv(rgba.zyxw, &h, &s, &v);

#else

		rgbToHsv(rgba, &h, &s, &v);

#endif

		uiHsvBuf[uiHsvBufIdx] = hsvPackToUint(h, s, v);

	}

}
3

I’m having the same problem, but it is not related to image2d_t. I switched to developer driver 260.61 yesterday (on Windows XP).

The code below will crash clBuildProgram() with an access violation. Comment out the kernel’s body (leaving an empty kernel, i.e. {}) and it will still crash, comment out the remaining functions and it will build.

I have two or three different kernels crashing that way.

void rgbToHsv(uint4 rgb, uint* uiH, uint* uiS, uint* uiV)

{  

	float r = rgb.x;

	float g = rgb.y;

	float b = rgb.z;

	float max = max(r, max(g, b));

	float min = min(r, min(g, b));

	float delta = max - min;

	*uiV = (uint) max;

	if (max == 0.0f) // r = g = b = 0

	{

		*uiS = 0;

		*uiH = 0;

		return;

	}

	else 

		*uiS = convert_uint_rte(255.0f * delta / max);

	float fH = 0.0f; // Case: max == min

	

	if (max != min)

	{

		float fInvDelta = 1.0f / delta;

		if (r == max)

			fH = (g - b) * fInvDelta;

		else if (g == max)

			fH = 2.0f + (b - r) * fInvDelta;

		else

			fH = 4.0f + (r - g) * fInvDelta;

	}

	fH *= 60.0f;

		

	if (fH < 0.0f)

		fH += 360.0f;

	

	*uiH = convert_uint_rte(fH);

}

uint hsvPackToUint(uint h, uint s, uint v)

{

	uint i = 0x000000FF & v;

	i |= 0x0000FF00 & (s << 8);

	i |= 0xFFFF0000 & (h << 16);

	return i;

}

__kernel void 

RgbBufToHsvBuf(

	__global __read_only uchar* ucRgbBuf,

	__global __write_only uint* uiHsvBuf,

	uint uiWidth,

	uint uiHeight)

{

	uint globalX = get_global_id(0);

	

	if (globalX >= uiWidth)

	{

		return;

	}

	for (uint y = 0, uiHsvBufIdx = globalX;

		y != uiHeight;

		++y, uiHsvBufIdx += uiWidth)

	{

		uint4 rgba;

		uint offset = uiHsvBufIdx * 3;

		rgba.x = ucRgbBuf[offset];

		rgba.y = ucRgbBuf[offset + 1];

		rgba.z = ucRgbBuf[offset + 2];

		rgba.w = 255;

	

		uint h, s, v;

#ifdef RGB_REVERSED

		rgbToHsv(rgba.zyxw, &h, &s, &v);

#else

		rgbToHsv(rgba, &h, &s, &v);

#endif

		uiHsvBuf[uiHsvBufIdx] = hsvPackToUint(h, s, v);

	}

}
4

Two things, first thank you SOOOOO much for solving my problem. You have to actually use the image2d_t, otherwise it crashes. Second, yes, I can verify that this is a problem, under a slightly different environment: Ubuntu 10.10 Linux Kernel 2.6.35-23-generic, driver version NVIDIA 260.19.14.

The behavior was exactly the same: if an image2d_t appears as a kernel argument, it MUST be used at least once in the kernel, otherwise, it will cause a segmentation fault when clBuildProgram is called. This is a bug because clBuildProgram must either succeed or let us print the error log. It cannot seg fault (or access violation) if the kernel source code looks wrong. That makes it extremely hard for developers to pin down the problem (it already wasted six hours of my time). I am not sure if the bug report has been filed, but I am putting a report together.

Moral of the story is it works fine, just make sure you use both the sampler and the image2d_t at least once.

5

Two things, first thank you SOOOOO much for solving my problem. You have to actually use the image2d_t, otherwise it crashes. Second, yes, I can verify that this is a problem, under a slightly different environment: Ubuntu 10.10 Linux Kernel 2.6.35-23-generic, driver version NVIDIA 260.19.14.

The behavior was exactly the same: if an image2d_t appears as a kernel argument, it MUST be used at least once in the kernel, otherwise, it will cause a segmentation fault when clBuildProgram is called. This is a bug because clBuildProgram must either succeed or let us print the error log. It cannot seg fault (or access violation) if the kernel source code looks wrong. That makes it extremely hard for developers to pin down the problem (it already wasted six hours of my time). I am not sure if the bug report has been filed, but I am putting a report together.

Moral of the story is it works fine, just make sure you use both the sampler and the image2d_t at least once.

6

Confirm also on platform:
Linux mylly 2.6.32-25-generic #45-Ubuntu SMP Sat Oct 16 19:52:42 UTC 2010 x86_64 GNU/Linux

generates dmesg:
[ 2474.222306] xxx[1707]: segfault at 3c ip 00007f161966c2de sp 00007ffffbbf6dc0 error 4 in libcuda.so.260.19.14[7f16195c8000+783000]

7

Confirm also on platform:
Linux mylly 2.6.32-25-generic #45-Ubuntu SMP Sat Oct 16 19:52:42 UTC 2010 x86_64 GNU/Linux

generates dmesg:
[ 2474.222306] xxx[1707]: segfault at 3c ip 00007f161966c2de sp 00007ffffbbf6dc0 error 4 in libcuda.so.260.19.14[7f16195c8000+783000]