device function wont return - gives unknown error will return when I hardcode a float but not from a

I am hoping for some help with a problem I am having with a device function. This function is called from my kernel. I want the function to interpolate a value for me. When I return the calculated value the function doesn’t work and gives “unknown error”. When I hard code a return value it works well. Please offer any suggestions you may have as I have exhausted the resources allocated by my brain at the moment. Thanks!

Here is the device function called from a working kernel:

__device__ float g_GetPartDepthAt(float3* pathPosition, float4* d_partTopSurface)


	float voxelSize = 0.08025;

 	int xDim = 104;


	int xCeil = ceilf( (pathPosition->x) / voxelSize);

	int xFloor = floorf( (pathPosition->x) / voxelSize);

	float xDelta = (pathPosition->x) / voxelSize - xFloor;


	int yCeil = ceilf( (pathPosition->y) / voxelSize);

	int yFloor = floorf( (pathPosition->y) / voxelSize);

	float yDelta = (pathPosition->y) / voxelSize - yFloor;

	int bottLeft = yFloor * xDim + xFloor;

	int topRight = yCeil * xDim + xCeil;

	int bottRight = yFloor * xDim + xCeil;

	int topLeft = yCeil * xDim + xFloor;

	float distX0_Y0 = d_blankTopSurface[bottLeft].w;

	float distX1_Y1 = d_blankTopSurface[topRight].w;

	float distX1_Y0 = d_blankTopSurface[bottRight].w;

	float distX0_Y1 = d_blankTopSurface[topLeft].w;

	float value1 = distX0_Y0 + xDelta * (distX1_Y0 - distX0_Y0);

	float value2 = distX0_Y1 + xDelta * (distX1_Y1 - distX0_Y1);

	float ans = (value1 + yDelta * (value2 - value1)); 

	return 1.0;

	//return ans;	


In the above function, I would like to return ans. Unfortunately, this doesn’t work. It works when I return a hard coded value like 1.0 as shown.

Any help you can offer is greatly appreciated.



Either [font=“Courier New”]pathPosition[/font] or [font=“Courier New”]d_blankTopSurface[/font] have bogus values, or you produce an out-of-bounds access to [font=“Courier New”]d_blankTopSurface[/font].

If you return a hard-coded value, the compiler will optimize out the memory accesses.

Is it intentional that [font=“Courier New”]d_partTopSurface[/font] is not used anywhere?


Thanks for your reply. It was not intentional meant for that to be


instead of


and this is why I was getting the problem. The reason, it was going as far as it did was because both variables are defined within the global scope. Sometimes you just need fresh eyes…