 texture<Scalar4, 1, cudaReadModeElementType> rigid_data_force_tex;
 //! The texture for reading the rigid data torque array
 texture<Scalar4, 1, cudaReadModeElementType> rigid_data_torque_tex;
-
 //! The texture for reading the net force array
 texture<Scalar4, 1, cudaReadModeElementType> net_force_tex;
+#elif defined ENABLE_TEXTURES
+//! The texture for reading the rigid data vel array
+texture<int4, 1, cudaReadModeElementType> rigid_data_vel_tex;
+//! The texture for reading the rigid data angvel array
+texture<int4, 1, cudaReadModeElementType> rigid_data_angvel_tex;
+//! The texture for reading the rigid data angmom array
+texture<int4, 1, cudaReadModeElementType> rigid_data_angmom_tex;
+//! The texture for reading the rigid data force array
+texture<int4, 1, cudaReadModeElementType> rigid_data_force_tex;
+//! The texture for reading the rigid data torque array
+texture<int4, 1, cudaReadModeElementType> rigid_data_torque_tex;
+//! The texture for reading the net force array
+texture<int4, 1, cudaReadModeElementType> net_force_tex;
+
+//! fetch_double4 Function for fetching double4 values from int4 textures
+/*! This function is only used when hoomd is compiled for double precision on the GPU.
+	
+	\param double_tex Texture in which the values are stored.
+	\param ii Index of the particle to read
+*/
+static __device__ inline Scalar4 fetch_double4(texture<int4, 1> double_tex, int ii)
+{
+	int idx = 2*ii;
+	int4 part1 = tex1Dfetch(double_tex, idx);
+	int4 part2 = tex1Dfetch(double_tex, idx+1);
+
+	return make_scalar4(__hiloint2double(part1.y, part1.x),
+		__hiloint2double(part1.w, part1.z),
+		__hiloint2double(part2.y, part2.x),
+		__hiloint2double(part2.w, part2.z));
+}
 #endif
 
 //! Shared memory used in reducing sums

				#ifdef SINGLE_PRECISION
                force = tex1Dfetch(rigid_data_force_tex, idx_body);
                vel = tex1Dfetch(rigid_data_vel_tex, idx_body);
				#else
				force = rdata_force[idx_body];
				vel = rdata_vel[idx_body];

    dim3 grid(num_blocks, 1, 1);
    dim3 threads(block_size, 1, 1);
    
    #ifdef ENABLE_TEXTURES
	cudaError_t error = cudaBindTexture(0, rigid_data_body_indices_tex, rdata.body_indices, sizeof(Scalar) * n_group_bodies);
    if (error != cudaSuccess)
        return error;
