atomicCAS() does not change the value (when executed with LLVM)

[GoogleCodeExporter]

What steps will reproduce the problem?
1. Run the following kernel, e.g. with 1 thread and 1 block:
__global__ void atomicCASTest(unsigned int* x, unsigned int* ret){
    unsigned int n = 2;
    *x = n;
    unsigned int x_old = atomicCAS(x, n, 5);
    ret[0] = n;
    ret[1] = x_old;
    ret[2] = *x;
}

2. Check the results of array "ret".

What is the expected output? What do you see instead?
Expected output is ret[0]=2, ret[1]=2, ret[2]=5. This is the case when executed 
with emulator (configured in configure.ocelot). 
However when executed with llvm, the result is ret[0]=2, ret[1]=2, ret[2]=2, so 
atomicCAS() did not change the value of x. 

What version of the product are you using? On what operating system?
Using llvm 3.2-6 and gpuocelot-svn built on 2013-05-01, on Arch Linux.

Please provide any additional information below.
Previously gpuocelot-svn on Arch Linux only compiled with llvm-svn. Therefore I 
patched the gpuocelot-svn build script to work with stable llvm version and to 
automatically download and build against LLVM header files. Maybe I made some 
mistake in the build script, so for reference here is the script:
https://aur.archlinux.org/packages/gp/gpuocelot-svn/PKGBUILD

Original issue reported on code.google.com by max.m...@dameweb.de on 11 May 2013 at 9:04

[GoogleCodeExporter]

Now I have tried it with llvm-svn and llvm 3.1 (and for each a rebuild of 
gpuocelot) as well but the result was the same. 
Maybe there is some error in gpuocelot when forwarding the atomic function to 
llvm?

Original comment by max.m...@dameweb.de on 12 May 2013 at 12:42

• Added labels: ****
• Removed labels: ****