         unsigned int idx = d_group_members[group_idx];
    
         // update positions to the next timestep and update velocities to the next half step
-        float4 pos = tex1Dfetch(pdata_pos_tex, idx);
+        float4 pos = pdata.pos[idx];
         
         float px = pos.x;
         float py = pos.y;
         float pz = pos.z;
         float pw = pos.w;
         
-        float4 vel = tex1Dfetch(pdata_vel_tex, idx);
-        float4 accel = tex1Dfetch(pdata_accel_tex, idx);
+        float4 vel = pdata.vel[idx];
+        float4 accel = pdata.accel[idx];
         
         vel.x = (vel.x + (1.0f/2.0f) * accel.x * deltaT) * denominv;
         px += vel.x * deltaT;

    dim3 grid( (group_size/block_size) + 1, 1, 1);
    dim3 threads(block_size, 1, 1);
    
    // bind the textures
    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_nvt_step_two_kernel<<< grid, threads >>>(pdata, d_group_members, group_size, Xi, deltaT);
    
    return cudaSuccess;
    }

void gpu_nvt_step_two_kernel(gpu_pdata_arrays pdata,
                             unsigned int *d_group_members,
                             unsigned int group_size,
                             float4 *d_net_force,
                             float Xi,
                             float deltaT)
    {
