3D vector CUDA kernel
The indentation within update() is misleading.  Most of all, everything inside of the function should be indented.  Otherwise, it'll be hard to tell what exactly is inside of the function.

I do, however, like that you've indented the (lengthy) parameter list by eight spaces instead of the standard four.  It should be easier to tell them apart from the code inside of the function.
More whitespace within statements can be added for more readability.

This line, for instance, is hard to read:


Ap[i]=cX*(p[i-sX]+p[i+sX]-2.0f*p[i])



You can add more whitespace as such:

Ap[i] = cX * (p[i-sX] + p[i+sX] - 2.0f * p[i])


Now you can really tell what everything is at a glance.

You can also add a space before the opening curly braces for functions.  It may not be a big deal with functions with smaller signatures, but yours are pretty lengthy, and it's harder to see the opening curly brace with them.

You also have inconsistent whitespace, such as in here:


for(int offset = blockDim.x/2;offset&gt;0;offset&gt;&gt;=1){



There is spacing between the =, but nowhere else.  This should look clearer:

for (int offset = blockDim.x / 2; offset &gt; 0; offset &gt;&gt;= 1) {

Considering that CUDA's blockIdx and blockDim are primarily signed, you should make your thread ID variables of the same type.  If they won't be too large, make them int.  If they can be larger than an int, make them either long or long long and cast blockIdx to the same.