binary search

Author: ngc92

src/models/llama_model.cpp

@@ -384,23 +384,18 @@ void LLamaModel::backward(Tensor inputs, Tensor targets, NCCLCommunicator& comm,
     rmsnorm_backward(rs->DActs[L-1].DResFFN.Value, d_lnf_w, rs->RMSNormScratch, rs->DActs[L - 1].DResFFN.Value, rs->DLNF,
                      rs->get_res_ffn(L-1, main_stream), Parameters->get_lnf(main_stream), rs->LNF_Rstd,
                      quant_abs_max_ptr(rs->DActs[L-1].DResFFN), B, T, C, rs->DeviceProp, main_stream);
-    CUDA_CHECK(cudaDeviceSynchronize());
     rs->release_res_ffn(L-1, main_stream);
 
     Parameters->release_lnf(main_stream);
-    CUDA_CHECK(cudaDeviceSynchronize());
     Grads->notify_lnf_w(main_stream, comm);
-    CUDA_CHECK(cudaDeviceSynchronize());
     rs->fetch_res_ffn(L-2, comm.stream());
-    CUDA_CHECK(cudaDeviceSynchronize());
     Parameters->gather_block(L - 1, comm, *rs);
     // now backward all the layers
     for (int l = L-1; l >= 0; l--) {
         NvtxRange layer_range("Layer", l);
         auto& dw = Grads->get_block_full(l, main_stream, comm, accumulate);
 
         // prefetch previous layer
-        CUDA_CHECK(cudaDeviceSynchronize());
         if(l > 1) {
             rs->fetch_res_ffn(l-2, comm.stream());
         }