cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

rick_weber
Adept II

get_image_dim returns wrong result on radeon 5870

When I run my kernel on the CPU, get_image_dim() returns the correct result, but the wrong result when run on my Radeon 5870. Also, what precisely are the restrictions on image dimensions? The width must be a multiple of 64? Height and width must be a power of 2?

0 Likes
5 Replies

There should be no restrictions at the OpenCL level on images. We are passing OpenCL conformance which tests this function directly so you might have hit a corner case. As always a test case that shows the problem would be greatly beneficial in helping us track it down.
0 Likes

This problem is really bizarre. It happens under the attached example. If I change #if 1 to #if 0, the problem goes away. Alternatively, I can get rid of types.cl and not pass the 2nd parameter. All other code is commented out. To make matters even weirder, only the dimesions on database are wrong, while the dimensions on spectrum are always correct.

When I allocated database, I got CL_SUCCESS, and have run no kernels between allocation and this. I have run other allocations and deviceWrites. I'm going to remove even more stuff and see if I can further isolate the problem. Furthermore, this problem occurs ONLY on my Radeon 5870; the results are correct on my Core i7.

main.cl: #include <specmaster.cl> #include <debug.cl> types.cl: #pragma once typedef struct _MSScanInfo { unsigned int mScanNum; unsigned int mMinMZ; unsigned int mMaxMZ; unsigned int mTotalCurrent; unsigned int mMSLevel; unsigned int mNumPeaks; unsigned int mPrecursorMZ; unsigned int mPrecursorCharge; unsigned int mNumHoles; unsigned int mCharge; unsigned int mIsGoodSpectra; }MSScanInfo; typedef enum _BYIon { kBIon, kYIon }BYIon; specmaster.cl: #pragma once //#include <clUtilImage1D.cl> //#include <tables.cl> #include <types.cl> //#include <scoreHelpers.cl> //#include <constants.cl> #pragma OPENCL EXTENSION cl_amd_printf : enable #if 0 #if 0 #define __BARRIER(type) \ printf("\t__BARRIER hit by thread %d\n", get_local_id(0));\ barrier(type); #else #define __BARRIER(type) \ barrier(type); #endif #endif #if 1 __kernel void foo(write_only image2d_t bar) { int2 coord = (get_global_id(0), get_global_id(1)); float4 output = (0.0f, 0.0f, 0.0f, 0.0f); write_imagef(bar, coord, output); } #endif __kernel void matchSpectra(read_only image2d_t spectrum, __global MSScanInfo* scanInfo, read_only image2d_t database, __global unsigned int* peptideMasses, unsigned int numPeptides, __global float* peptideScores, __global unsigned int* peptideScoresIdx) { peptideScoresIdx[0] = get_image_width(database); peptideScoresIdx[1] = get_image_width(spectrum); } debug.cl: __kernel void viewMemory(__global void* memory) { return; } __kernel void getImageInfo(read_only image2d_t theImage, __global unsigned int* info) { info[0] = get_image_width(theImage); info[1] = get_image_height(theImage); }

0 Likes

To make matters stranger, getImageInfo DOES return the correct the correct dimensions for the image, and I call it immediately after matchSpectra(). This means the image isn't corrupted on the device.

I've commented out pretty much everything except allocations used by the matchSpectra() kernel, so there's nothing else that can be causing this. I can access the pixels themselves, but when the dimensions returned are incorrect, my program can't emulate a 1D texture like it needs. I guess there are two workarounds: explicitly pass the leading dimension to the kernel or get rid of other routines that access images in the same file. Neither is attractive, since the former since is cumbersome and the latter requires changing a library I have.

0 Likes

Can you provide a proper test case with runtime code. It would be very easy to reproduce the issue at our end.

0 Likes

Other than the source code I pasted, I don't know if I could construct a proper testcase, because it seems to come up only when the price of tea in China exceeds $20 and ounce.

Revision 68 of

http://code.google.com/p/specmaster/source/checkout

was the last to contain the bug. r69 fixed it using a workaround previously described.

To build, you'll need libXML, clUtil, and g++-4.5. libXML is in package libxml-dev.

I'm not sure if gcc 4.5 is in Ubuntu 10.04's repositories yet - I had to build it from scratch. I know it's in apt-get in 11.04 as g++-4.5, but that's not officially supported (though, from experience it seems to work).

You can get clUtil at:

http://code.google.com/p/clutil/source/checkout

If it's any more complicated than typing make, let me know.

 

0 Likes