 __host__ void ChSystemGpu_impl::runSphereBroadphase() {
     METRICS_PRINTF("Resetting broadphase info!\n");
 
-    // reset the number of spheres per SD, the offsets in the big composite array, the big fat composite array, and the
-    // scrath pad
+    // reset the number of spheres per SD, the offsets in the big composite array, and the big fat composite array
     resetBroadphaseInformation();
 
     // Frist stage of the computation in this function: Figure out the how many spheres touch each SD. 

    unsigned int* in_ptr = SD_NumSpheresTouching.data();
    gpuErrchk(cudaMemcpy(out_ptr, in_ptr, nSDs * sizeof(unsigned int), cudaMemcpyDeviceToDevice));

    // cold run; CUB determines the amount of storage it needs (since d_temp_storage is NULL)
    void* d_temp_storage = NULL;
    size_t temp_storage_bytes = 0;
    cub::DeviceScan::ExclusiveSum(d_temp_storage, temp_storage_bytes, in_ptr, out_ptr, nSDs);
    gpuErrchk(cudaDeviceSynchronize());
    gpuErrchk(cudaPeekAtLastError());
   
    // provide CUB with needed temporary storage
    if (d_temp_storageCUB.size() < temp_storage_bytes) 
        d_temp_storageCUB.resize(temp_storage_bytes);
    sphere_data->d_temp_storageCUB = d_temp_storageCUB.data();
    // Run the actual exclusive prefix sum
    cub::DeviceScan::ExclusiveSum(sphere_data->d_temp_storageCUB, temp_storage_bytes, in_ptr, out_ptr, nSDs);
    gpuErrchk(cudaDeviceSynchronize());
    gpuErrchk(cudaPeekAtLastError());


    // Run the simulation, there are aggressive synchronizations because we want to have no race conditions
    for (unsigned int n = 0; n < nsteps; n++) {
        updateBCPositions();
        runSphereBroadphase();
        packSphereDataPointers();
        resetSphereAccelerations();
        resetBCForces();

