how do I gather() a variable in atomicAdd()?

I am just wondering how to return result from atomicAdd(), since atomicAdd only changes one of its inputs. Here is my kernel (CUDA C):

__global__ void agpuhelloworld( double* c, double ato, double b ) {
*c = ato + b;
atomicAdd(&ato, b);

}

I know how to collect c, but how do I collect value of ato from this kernel with a MATLAB script? I cannot declare that as persistent, because ato is one of the inputs to MATLAB function. Following is my MATLAB code, I think it is something like that. I only show how to gather() c.

function out = gpuhelloworld(ato,b)
persistent gpuadd;
persistent c;

if isempty(gpuadd)
    disp 'Initializing GPU drf calculation...';
    gpuadd = parallel.gpu.CUDAKernel('agpuhelloworld.ptx','agpuhelloworld.cu','agpuhelloworld');
    tmp = gpuDevice;
    gpuadd.GridSize = [min(1024,tmp.MaxGridSize(1)) 1]; % min(1024, maxblocks)
    gpuadd.ThreadBlockSize = [min(128,tmp.MaxThreadsPerBlock) 1 1]; % min(128, maxthreadsperblock)
    c = gpuArray(zeros(1,1));

end;

c= feval(gpuadd, c,ato,b);
out = gather(c);

Answer 1

[m4 Chrennikov]

Try this:

[c, ato] = feval(gpuadd, c,ato,b);
out = gather(ato);

Answer 2

[Nick Tsui]

Thanks it works.