     if (idx < pdata.N)
         {
         // fetch particle position
-        float4 pos = tex1Dfetch(pdata_pos_tex, idx);
+        float4 pos = pdata.pos[idx];
 
         float px = pos.x;
         float py = pos.y;

        float pw = pos.w;

        // read in the image flags
        int4 image = tex1Dfetch(pdata_image_tex, idx);

        // fix periodic boundary conditions
        float x_shift = rintf(px * box.Lxinv);

    unsigned int block_size = 256;
    dim3 grid( (group_size / block_size) + 1, 1, 1);
    dim3 threads(block_size, 1, 1);

    // run the kernel
    gpu_nph_step_one_kernel<<< grid, threads >>>(pdata,
