         return error;
     
     // run the kernel
-    gpu_bdnvt_step_two_kernel<<< grid, threads >>>(pdata,
+    gpu_bdnvt_step_two_kernel<<< grid, threads, sizeof(float)*bdnvt_args.n_types >>>
+                                                  (pdata,
                                                    d_group_members,
                                                    group_size,
                                                    d_net_force,

        // ******** first, calculate the additional BD force
        // read the current particle velocity (MEM TRANSFER: 16 bytes)
        float4 vel = tex1Dfetch(pdata_vel_tex, idx);
        // read in the type of our particle. A texture read of only the fourth part of the position float4 (where type is stored) is used.
        // (MEM TRANSFER: 4 bytes)
        unsigned int typ = tex1Dfetch(pdata_type_tex, idx*4 + 3);
        // read in the tag of our particle.
        // (MEM TRANSFER: 4 bytes)
        unsigned int ptag = tex1Dfetch(pdata_tag_tex, idx);

            }
        else
            {
            // read in the type of our particle. A texture read of only the fourth part of the position float4
            // (where type is stored) is used.
            unsigned int typ = tex1Dfetch(pdata_type_tex, idx*4 + 3);
            gamma = s_gammas[typ];
            }
        
