// CPU time
  struct timeval tp;
  gettimeofday(&tp, NULL);
  double iStart = ((double)tp.tv_sec + (double)tp.tv_usec * 1.e-6);

  SparseFillEmptyRows(input_indice_addr, input_values_addr, input_default_values_addr, input_dense_shape_addr, 0,
                      input_indices_shapes_[0], dense_row, workspace_elements_per_rows_addr,
                      workspace_empty_rows_count_addr, workspace_row_indices_addr, workspace_input_row_ends_addr,
                      workspace_sorted_indices_addr, workspace_final_shape_addr, workspace_origin_index_addr,
                      workspace_sorted_key_addr, reinterpret_cast<cudaStream_t>(cuda_stream_), output_indices_addr,
                      output_values_addr, output_empty_row_indicator_addr, output_reverse_index_map_addr);
  CHECK_CUDA_RET_WITH_EXCEPT_NOTRACE(
    cudaMemcpyAsync(&real_output_size_, workspace_final_shape_addr, sizeof(int64_t), cudaMemcpyDeviceToHost,
                    reinterpret_cast<cudaStream_t>(cuda_stream_)),
    "SparseFillEmptyRows cudaMemcpyAsync failed.");
  cudaDeviceSynchronize();  // 同步函数
  gettimeofday(&tp, NULL);
  double iElaps = ((double)tp.tv_sec + (double)tp.tv_usec * 1.e-6) - iStart;  //单位是second
  std::cout << "CPU TIME 237-246:" << iElaps * 1.e3 << std::endl;
  // CPU time end
  return true;

  cudaEvent_t start, stop;
  float elapsedTime = 0.0;
  cudaEventCreate(&start);
  cudaEventCreate(&stop);
  cudaEventRecord(start, 0);

  CopyRowIndiceKernel<<<CUDA_BLOCKS(device_id, indice_num), GET_THREADS, 0, cuda_stream>>>(indices_ptr, row_indices,
                                                                                           origin_index, indice_num);
  cudaEventRecord(stop, 0);
  cudaMemset(elements_per_rows, 0, dense_row * sizeof(int64_t));
  CalElementPerRowsKernel<<<CUDA_BLOCKS(device_id, indice_num), GET_THREADS, 0, cuda_stream>>>(
    dense_shape_ptr, row_indices, elements_per_rows, indice_num);
  CalEmptyRowIndicatorKernel<<<CUDA_BLOCKS(device_id, dense_row), GET_THREADS, 0, cuda_stream>>>(
    elements_per_rows, dense_row, output_empty_row_indicator_ptr);
  InclusivePrefixSum(dense_row, elements_per_rows, input_row_ends, cuda_stream);
  InclusiveBoolPrefixSum(dense_row, output_empty_row_indicator_ptr, empty_row_count_sum, cuda_stream);
  RowsSort(indice_num, cuda_stream, dense_shape_ptr, row_indices, origin_index, sorted_key, sorted_indices, device_id);
  AssignValueKernel<<<CUDA_BLOCKS(device_id, dense_row), GET_THREADS, 0, cuda_stream>>>(
    values_ptr, indices_ptr, sorted_indices, dense_row, default_value, empty_row_count_sum, input_row_ends,
    output_values_ptr, output_indices_ptr, indice_num, final_shape, output_reverse_index_map_ptr);

  cudaEventRecord(stop, 0);
  cudaEventSynchronize(stop);
  cudaEventElapsedTime(&elapsedTime, start, stop);
  std::cout << "AssignValueKernel" << elapsedTime << std::endl;
  cudaEventDestroy(start);
  cudaEventDestroy(stop);
  std::cout << std::endl;

 # get random data.
def gen_data_file(shape, dtype, rand_type, low, high):
   if rand_type == "randint":
       rand_data = np.random.randint(low, high, size=shape)
   else:
       rand_data = np.random.uniform(low, high, size=shape)
   data = np.array(rand_data, dtype=dtype)
   return data


# generate dense_shape, limit_cnt
def get_rand_dense():
   dense_x = random.randint(1, 1000)
   dense_y = random.randint(1, 1000)
   limit = dense_x*dense_y
   return dense_x, dense_y, limit

# 生成规定cnt数目， 在limit_x，limit_y范围内的indices 并且不重复。
def rand_index(cnt, limit_x, limit_y):
   res = set()
   for i in range(cnt):
       while True:
           x = random.randint(0, limit_x - 1)
           y = random.randint(0, limit_y - 1)
           if (x, y) in res:
               continue
           res.add((x, y))
           break
   return Tensor(list(res))
'''
   dense 2000W, values number=10W
'''
def test_sfer_2000x10000_performance():
   indices = rand_index(100000, 10000, 2000)
   values =Tensor(np.array(np.random.randn(100000)),dtype=ms.float32)
   default_value = Tensor(np.random.randn(),dtype = ms.float32)
   dense = Tensor([10000, 2000], dtype=ms.int64)
   fact = SparseFillEmptyRowsMock(inputs=[indices, values, dense, default_value])
   fact.forward_profile_cmp()