         unsigned int idx = d_group_members[group_idx];
 
         // fetch particle position
-        float4 pos = tex1Dfetch(pdata_pos_tex, idx);
+        float4 pos = pdata.pos[idx];
 
         float px = pos.x;
         float py = pos.y;

    dim3 grid( (group_size / block_size) + 1, 1, 1);
    dim3 threads(block_size, 1, 1);

    // bind the texture
    cudaError_t error = cudaBindTexture(0, pdata_vel_tex, pdata.vel, sizeof(float4) * pdata.N);
    if (error != cudaSuccess)
        return error;

    error = cudaBindTexture(0, pdata_mass_tex, pdata.mass, sizeof(float) * pdata.N);
    if (error != cudaSuccess)
        return error;

    error = cudaBindTexture(0, net_force_tex, d_net_force, sizeof(float4) * pdata.N);
    if (error != cudaSuccess)
        return error;

    // run the kernel
    gpu_nph_step_two_kernel<<< grid, threads >>>(pdata, d_group_members, group_size, deltaT);

    return cudaSuccess;
    }

    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,
