cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

mjharvey
Adept I

beta4 problem: "link failed" when compiling a kernel using atomics

beta4 problem

Compiling the following kernel:

"#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable


__kernel void test( __global int *ptr ) {
        atom_add( ptr, 1 );
}

"

 

Gives the compiler error

"Link failed"

The failure is directly related to the use of atomic operations.

This is with the 64bit Ubuntu 9.04 release.

0 Likes
13 Replies
genaganna
Journeyman III

are you compiling for GPU or CPU?

 

It should work for CPU. It is expected for GPU as extenstions are not supported in beta4 release

0 Likes

Yes, GPU.

I'd missed that only the CPU target supports atomics. (As that's the case though, I would have expected the #pragma to cause the compiler to error out.)

 

 

 

0 Likes

Yes, you are right.

Compiler should give appropriate error.

0 Likes

mjharvey, what GPU are you trying to run this on? The 4XXX series of cards does not have atomic hardware and thus will never support atomic extensions.
0 Likes

On a related note: I just upgraded from beta2 to beta4, and I'm trying to build my program on the CPU to make sure it still works.  It uses atomics, so I'd expect it not to work on the GPU; however, clBuildProgram() returning an error based on the line:

    #pragma OPENCL EXTENSION cl_khr_local_in32_base_atomics : enable

The error it prints to the build log is: "warning: unrecognized STDC pragma".  And then all the atomic functions my kernel calls are reported as undefined.  I know the relevant atomic extension is supported; it's listed in the device extensions string.  Also, if I take the high road and enable all atomics:

    #pragma OPENCL EXTENSION all : enable

then everything works builds and runs correctly.  What could I be doing wrong?

UPDATE: this is under Windows XP 64, with a Radeon 4800-series GPU.

0 Likes

Micah,

 

That's interesting to learn that the 4xxx doesn't support atomic operations (I have a 4550 for testing purposes). Is there anywhere a comprehensive list of the features (un)supported by the 4xxx and 5xxx series cards?

Ta,

Matt

0 Likes

#pragma OPENCL EXTENSION cl_khr_local_in32_base_atomics : enable <-- typo, it is int32
0 Likes

...durp   Wonder how that ever worked, then.  Thanks!

0 Likes

I don't think we ever enforced it until Beta3/4
0 Likes

Other than looking at their specs?

4XXX series supports base 1.0 OpenCL Only, no extensions.
5XXX currently only supports 1.0 OpenCL base, but will support atomics, byte addressable in our next release.
0 Likes

what do you mean by byte addressable?

does 5870 has unified memory?

btw it seems that 5900 series are on the way, listed in sdk

0 Likes

Originally posted by: MicahVillmow Other than looking at their specs? 4XXX series supports base 1.0 OpenCL Only, no extensions. 5XXX currently only supports 1.0 OpenCL base, but will support atomics, byte addressable in our next release.


absolutly no extension? like double floating point?

 

riza.guntur: byte addressable see section 9.9 of OpenCL specification. normaly you cant write to array of data types les the 32-bits. with byte addressable it is possible.

0 Likes

Originally posted by: MicahVillmow Other than looking at their specs? 4XXX series supports base 1.0 OpenCL Only, no extensions. 5XXX currently only supports 1.0 OpenCL base, but will support atomics, byte addressable in our next release.


What! So OpenCL GPU for 4800 don't support double?

But double should work fine for CPU right? Or I'm I wrong?

0 Likes