     dim3 grid( (group_size/block_size) + 1, 1, 1);
     dim3 threads(block_size, 1, 1);
     
-    // bind the textures
-    cudaError_t error = cudaBindTexture(0, pdata_pos_tex, pdata.pos, sizeof(float4) * pdata.N);
-    if (error != cudaSuccess)
-        return error;
-        
-    error = cudaBindTexture(0, pdata_vel_tex, pdata.vel, sizeof(float4) * pdata.N);
-    if (error != cudaSuccess)
-        return error;
-        
-    error = cudaBindTexture(0, pdata_accel_tex, pdata.accel, sizeof(float4) * pdata.N);
-    if (error != cudaSuccess)
-        return error;
-        
-    error = cudaBindTexture(0, pdata_image_tex, pdata.image, sizeof(int4) * pdata.N);
-    if (error != cudaSuccess)
-        return error;
-
-        
     // run the kernel
     gpu_nvt_step_one_kernel<<< grid, threads, block_size * sizeof(float) >>>(pdata,
                                                                              d_group_members,

#include <assert.h>
#endif

//! The texture for reading the pdata pos array
texture<float4, 1, cudaReadModeElementType> pdata_pos_tex;
//! The texture for reading the pdata vel array
texture<float4, 1, cudaReadModeElementType> pdata_vel_tex;
//! The texture for reading the pdata accel array
texture<float4, 1, cudaReadModeElementType> pdata_accel_tex;
//! The texture for reading in the pdata image array
texture<int4, 1, cudaReadModeElementType> pdata_image_tex;
//! The texture for reading particle mass
texture<float, 1, cudaReadModeElementType> pdata_mass_tex;

//! Shared memory used in reducing the 2K sum
extern __shared__ float nvt_sdata[];

/*! \file TwoStepNVTGPU.cu
    \brief Defines GPU kernel code for NVT integration on the GPU. Used by TwoStepNVTGPU.
*/

        unsigned int idx = d_group_members[group_idx];
   
        // read in the net force and calculate the acceleration
        float4 accel = d_net_force[idx];
        float mass = pdata.mass[idx];
        accel.x /= mass;
        accel.y /= mass;
        accel.z /= mass;
        
        float4 vel = pdata.vel[idx];
        
        vel.x += (1.0f/2.0f) * deltaT * (accel.x - Xi * vel.x);
        vel.y += (1.0f/2.0f) * deltaT * (accel.y - Xi * vel.y);
