 //! Texture for reading particle positions
 texture<float, 1, cudaReadModeElementType> pdata_type_tex;
 
+//! Texture for reading particle tags
+texture<unsigned int, 1, cudaReadModeElementType> pdata_tag_tex;
+
 //! Kernel for calculating stochastic forces
 /*! This kernel is called to apply stochastic heat bath forces to all N particles in conjunction with a Brownian Dynamics Simulations
 

	float type_f = tex1Dfetch(pdata_type_tex, idx_global*4 + 3);
	int typ = __float_as_int(type_f);
	
	// Calculate Coefficient of Friction
	//type = 0;   //May use this for benchmarking the impact of doing a second texture read just for particle type
	float coeff_fric = sqrtf(6.0f * s_gammas[typ] * T/ dt);

	error = cudaBindTexture(0, pdata_type_tex, pdata->pos, sizeof(float4) * pdata->N);
	if (error != cudaSuccess)
		return error;
	
	// bind the tag texture
	error = cudaBindTexture(0, pdata_tag_tex, pdata->tag, sizeof(unsigned int) * pdata->N);
	if (error != cudaSuccess)
		return error;		
		
    // run the kernel
    stochasticForces_kernel<<< grid, threads, sizeof(float)*gamma_length>>>(force_data, *pdata, dt, T, d_gammas, gamma_length, seed, iteration);
