     dim3 grid( (pdata.N / 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_image_tex, pdata.image, sizeof(int4) * pdata.N);
-    if (error != cudaSuccess)
-        return error;
-
     // run the kernel
     gpu_nph_wrap_particles_kernel<<< grid, threads >>>(pdata, box);
 
     return cudaSuccess;
     }
-
-//! The texture for reading the net force
-texture<float4, 1, cudaReadModeElementType> net_force_tex;
 
 /*! \param pdata Particle data arrays to integrate forward 1/2 step
     \param d_group_members Device array listing the indicies of the mebers of the group to integrate

    if (idx < pdata.N)
        {
        // fetch particle position
        float4 pos = tex1Dfetch(pdata_pos_tex, idx);

        float px = pos.x;
        float py = pos.y;

    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,
