I am currently working my way through the CUDA By Example book, porting to PyCuda as I go. I've reached the first example that makes use of a struct. I'm a bit unsure of how to properly deal with this in PyCuda. I followed the documentation as best I could, and I wrote a kernel test out if my wrapper was working correctly. The kernel executes without error, but fails to alter data in the struct. Code follows, and is attached for easy viewing in your favorite editor.

Best,
Nathan

import pycuda.driver as cuda
from pycuda.compiler import SourceModule
import pycuda.autoinit
import numpy

if __name__ == "__main__":
    mod = SourceModule('''
        // Code taken from CUDA by Example
        
        #define INF 2e10f
        
        struct Sphere {
            float x, y, z;
            float radius;
            float r, g, b;
            
            __device__ float hit(float ox, float oy, float* n){
                float dx = ox - x;
                float dy = oy - y;
                if (dx*dx + dy*dy < radius*radius) {
                    float dz = sqrtf(radius*radius - dx*dx - dy*dy);
                    *n = dz / sqrtf(radius*radius);
                    return dz + z;
                }
                return -INF;
            }
        };
        
        // Test function I wrote for checking that my wrapper code works
        
        __global__ void addOne(Sphere* s){
            s = &s[threadIdx.x];
            s->x += 1.0f;
            s->y += 1.0f;
            s->z += 1.0f;
            s->radius += 1.0f;
            s->r += 1.0f;
            s->g += 1.0f;
            s->b += 1.0f;
        }
    ''')

# wrapper code adapted from Nicholas Tung's example in the pycuda docs
class SphereStruct:
    # Sphere has 7 float fields (x, y, z, radius, r, g, b)
    mem_size = 7 * numpy.float32(0).nbytes
    def __init__(self, position, radius, color, struct_ptr):
        self.position = cuda.to_device(position)
        self.radius = cuda.to_device(radius)
        self.color = cuda.to_device(color)
        cuda.memcpy_htod(int(struct_ptr), position)
        cuda.memcpy_htod(int(struct_ptr) + numpy.float32(0).nbytes * (3), radius)
        cuda.memcpy_htod(int(struct_ptr) + numpy.float32(0).nbytes * (3+1), color)


    def __str__(self):
        return ("position: " + str(cuda.from_device(self.position, (3), numpy.float32)) + ", " +
                "radius: " + str(cuda.from_device(self.radius, (1), numpy.float32)) +
                " color: " + str(cuda.from_device(self.color, (3), numpy.float32)))

# allocate memory for the struct        
sphere_struct = cuda.mem_alloc(SphereStruct.mem_size)
# declare the the wrapper object
sphere = SphereStruct(numpy.array([1,2,3], dtype=numpy.float32),
                      numpy.float32(5),
                      numpy.array([255,100,20], dtype=numpy.float32),
                      sphere_struct)

print(str(sphere))
func = mod.get_function('addOne')
func(sphere_struct, grid = (1,1), block = (1,1,1))
print(str(sphere))