<commit_before>#include "lm_fit.h"
#include <algorithm>
#include "cuda_kernels.cuh"

#ifdef ARCH_64

void LMFitCUDA::solve_equation_systems_lup()
{
    dim3  threads(1, 1, 1);
    dim3  blocks(1, 1, 1);

    // initialize components of equation systems
    gpu_data_.copy(gpu_data_.decomposed_hessians_, gpu_data_.hessians_, n_fits_ * info_.n_parameters_to_fit_ * info_.n_parameters_to_fit_);

    // decompose hessians
    cublasStatus_t lu_status_decopmposition = cublasSgetrfBatched(
        gpu_data_.cublas_handle_,
        info_.n_parameters_to_fit_,
        gpu_data_.pointer_decomposed_hessians_,
        info_.n_parameters_to_fit_,
        gpu_data_.pivot_vectors_,
        gpu_data_.solution_info_,
        n_fits_);

    // initialize deltas with values of gradients
    gpu_data_.copy(gpu_data_.deltas_, gpu_data_.gradients_, n_fits_ * info_.n_parameters_to_fit_);

    // TODO: check solution_info
    int solution_info;

    // solve equation systems
    cublasStatus_t lu_status_solution
        = cublasSgetrsBatched(
        gpu_data_.cublas_handle_,
        CUBLAS_OP_N,
        info_.n_parameters_to_fit_,
        1,
        (float const **)(gpu_data_.pointer_decomposed_hessians_.data()),
        info_.n_parameters_to_fit_,
        gpu_data_.pivot_vectors_,
        gpu_data_.pointer_deltas_,
        info_.n_parameters_to_fit_,
        &solution_info,
        n_fits_);
}

#else //ARCH_64

void LMFitCUDA::solve_equation_systems_gj()
{
    dim3  threads(1, 1, 1);
    dim3  blocks(1, 1, 1);

    int n_parameters_pow2 = 1;

    while (n_parameters_pow2 < info_.n_parameters_to_fit_)
    {
        n_parameters_pow2 *= 2;
    }

    //set up to run the Gauss Jordan elimination
    int const n_equations = info_.n_parameters_to_fit_;
    int const n_solutions = n_fits_;

    threads.x = n_equations + 1;
    threads.y = n_equations;
    blocks.x = n_solutions;

    //set the size of the shared memory area for each block
    int const shared_size
        = sizeof(float) * ((threads.x * threads.y)
            + n_parameters_pow2 + n_parameters_pow2);

    //run the Gauss Jordan elimination
    cuda_gaussjordan <<< blocks, threads, shared_size >>>(
        gpu_data_.deltas_,
        gpu_data_.gradients_,
        gpu_data_.hessians_,
        gpu_data_.finished_,
        gpu_data_.solution_info_,
        info_.n_parameters_to_fit_,
        n_parameters_pow2);
    CUDA_CHECK_STATUS(cudaGetLastError());
}

#endif // ARCH_64

void LMFitCUDA::update_states()
{
    dim3  threads(1, 1, 1);
    dim3  blocks(1, 1, 1);

    //set up to update the lm_state_gpu_ variable with the Gauss Jordan results
    threads.x = std::min(n_fits_, 256);
    blocks.x = int(std::ceil(float(n_fits_) / float(threads.x)));

    //update the gpu_data_.states_ variable
    cuda_update_state_after_solving <<< blocks, threads >>>(
        n_fits_,
        gpu_data_.solution_info_,
        gpu_data_.finished_,
        gpu_data_.states_);
    CUDA_CHECK_STATUS(cudaGetLastError());
}

void LMFitCUDA::scale_hessians()
{
    dim3  threads(1, 1, 1);
    dim3  blocks(1, 1, 1);

    threads.x = info_.n_parameters_to_fit_*info_.n_fits_per_block_;
    blocks.x = n_fits_ / info_.n_fits_per_block_;

    cuda_modify_step_widths <<< blocks, threads >>>(
        gpu_data_.hessians_,
        gpu_data_.lambdas_,
        gpu_data_.scaling_vectors_,
        info_.n_parameters_to_fit_,
        gpu_data_.iteration_failed_,
        gpu_data_.finished_,
        info_.n_fits_per_block_);
    CUDA_CHECK_STATUS(cudaGetLastError());
}

void LMFitCUDA::update_parameters()
{
    dim3  threads(1, 1, 1);
    dim3  blocks(1, 1, 1);

    threads.x = info_.n_parameters_*info_.n_fits_per_block_;
    blocks.x = n_fits_ / info_.n_fits_per_block_;

    cuda_update_parameters <<< blocks, threads >>>(
        gpu_data_.parameters_,
        gpu_data_.prev_parameters_,
        gpu_data_.deltas_,
        info_.n_parameters_to_fit_,
        gpu_data_.parameters_to_fit_indices_,
        gpu_data_.finished_,
        info_.n_fits_per_block_);
    CUDA_CHECK_STATUS(cudaGetLastError());
}

void LMFitCUDA::calc_curve_values()
{
	dim3  threads(1, 1, 1);
	dim3  blocks(1, 1, 1);

	threads.x = info_.n_points_ * info_.n_fits_per_block_ / info_.n_blocks_per_fit_;
    if (info_.n_blocks_per_fit_ > 1)
        threads.x += info_.n_points_ % threads.x;
	blocks.x = n_fits_ / info_.n_fits_per_block_ * info_.n_blocks_per_fit_;

	cuda_calc_curve_values <<< blocks, threads >>>(
		gpu_data_.parameters_,
		n_fits_,
		info_.n_points_,
		info_.n_parameters_,
		gpu_data_.finished_,
		gpu_data_.values_,
		gpu_data_.derivatives_,
		info_.n_fits_per_block_,
        info_.n_blocks_per_fit_,
		info_.model_id_,
		gpu_data_.chunk_index_,
		gpu_data_.user_info_,
		info_.user_info_size_);
	CUDA_CHECK_STATUS(cudaGetLastError());
}

void LMFitCUDA::calc_chi_squares()
{
    dim3  threads(1, 1, 1);
    dim3  blocks(1, 1, 1);

    threads.x = info_.power_of_two_n_points_ * info_.n_fits_per_block_ / info_.n_blocks_per_fit_;
    blocks.x = n_fits_ / info_.n_fits_per_block_ * info_.n_blocks_per_fit_;

    int const shared_size = sizeof(float) * threads.x;

    float * chi_squares = 
        info_.n_blocks_per_fit_ > 1 ? gpu_data_.subtotals_ : gpu_data_.chi_squares_;

    cuda_calculate_chi_squares <<< blocks, threads, shared_size >>>(
        chi_squares,
        gpu_data_.states_,
        gpu_data_.data_,
        gpu_data_.values_,
        gpu_data_.weights_,
        info_.n_points_,
        n_fits_,
        info_.estimator_id_,
        gpu_data_.finished_,
        info_.n_fits_per_block_,<commit_msg>fixed whitespace<commit_after>#include "lm_fit.h"
#include <algorithm>
#include "cuda_kernels.cuh"

#ifdef ARCH_64

void LMFitCUDA::solve_equation_systems_lup()
{
    dim3  threads(1, 1, 1);
    dim3  blocks(1, 1, 1);

    // initialize components of equation systems
    gpu_data_.copy(gpu_data_.decomposed_hessians_, gpu_data_.hessians_, n_fits_ * info_.n_parameters_to_fit_ * info_.n_parameters_to_fit_);

    // decompose hessians
    cublasStatus_t lu_status_decopmposition = cublasSgetrfBatched(
        gpu_data_.cublas_handle_,
        info_.n_parameters_to_fit_,
        gpu_data_.pointer_decomposed_hessians_,
        info_.n_parameters_to_fit_,
        gpu_data_.pivot_vectors_,
        gpu_data_.solution_info_,
        n_fits_);

    // initialize deltas with values of gradients
    gpu_data_.copy(gpu_data_.deltas_, gpu_data_.gradients_, n_fits_ * info_.n_parameters_to_fit_);

    // TODO: check solution_info
    int solution_info;

    // solve equation systems
    cublasStatus_t lu_status_solution
        = cublasSgetrsBatched(
        gpu_data_.cublas_handle_,
        CUBLAS_OP_N,
        info_.n_parameters_to_fit_,
        1,
        (float const **)(gpu_data_.pointer_decomposed_hessians_.data()),
        info_.n_parameters_to_fit_,
        gpu_data_.pivot_vectors_,
        gpu_data_.pointer_deltas_,
        info_.n_parameters_to_fit_,
        &solution_info,
        n_fits_);
}

#else //ARCH_64

void LMFitCUDA::solve_equation_systems_gj()
{
    dim3  threads(1, 1, 1);
    dim3  blocks(1, 1, 1);

    int n_parameters_pow2 = 1;

    while (n_parameters_pow2 < info_.n_parameters_to_fit_)
    {
        n_parameters_pow2 *= 2;
    }

    //set up to run the Gauss Jordan elimination
    int const n_equations = info_.n_parameters_to_fit_;
    int const n_solutions = n_fits_;

    threads.x = n_equations + 1;
    threads.y = n_equations;
    blocks.x = n_solutions;

    //set the size of the shared memory area for each block
    int const shared_size
        = sizeof(float) * ((threads.x * threads.y)
            + n_parameters_pow2 + n_parameters_pow2);

    //run the Gauss Jordan elimination
    cuda_gaussjordan <<< blocks, threads, shared_size >>>(
        gpu_data_.deltas_,
        gpu_data_.gradients_,
        gpu_data_.hessians_,
        gpu_data_.finished_,
        gpu_data_.solution_info_,
        info_.n_parameters_to_fit_,
        n_parameters_pow2);
    CUDA_CHECK_STATUS(cudaGetLastError());
}

#endif // ARCH_64

void LMFitCUDA::update_states()
{
    dim3  threads(1, 1, 1);
    dim3  blocks(1, 1, 1);

    //set up to update the lm_state_gpu_ variable with the Gauss Jordan results
    threads.x = std::min(n_fits_, 256);
    blocks.x = int(std::ceil(float(n_fits_) / float(threads.x)));

    //update the gpu_data_.states_ variable
    cuda_update_state_after_solving <<< blocks, threads >>>(
        n_fits_,
        gpu_data_.solution_info_,
        gpu_data_.finished_,
        gpu_data_.states_);
    CUDA_CHECK_STATUS(cudaGetLastError());
}

void LMFitCUDA::scale_hessians()
{
    dim3  threads(1, 1, 1);
    dim3  blocks(1, 1, 1);

    threads.x = info_.n_parameters_to_fit_*info_.n_fits_per_block_;
    blocks.x = n_fits_ / info_.n_fits_per_block_;

    cuda_modify_step_widths <<< blocks, threads >>>(
        gpu_data_.hessians_,
        gpu_data_.lambdas_,
        gpu_data_.scaling_vectors_,
        info_.n_parameters_to_fit_,
        gpu_data_.iteration_failed_,
        gpu_data_.finished_,
        info_.n_fits_per_block_);
    CUDA_CHECK_STATUS(cudaGetLastError());
}

void LMFitCUDA::update_parameters()
{
    dim3  threads(1, 1, 1);
    dim3  blocks(1, 1, 1);

    threads.x = info_.n_parameters_*info_.n_fits_per_block_;
    blocks.x = n_fits_ / info_.n_fits_per_block_;

    cuda_update_parameters <<< blocks, threads >>>(
        gpu_data_.parameters_,
        gpu_data_.prev_parameters_,
        gpu_data_.deltas_,
        info_.n_parameters_to_fit_,
        gpu_data_.parameters_to_fit_indices_,
        gpu_data_.finished_,
        info_.n_fits_per_block_);
    CUDA_CHECK_STATUS(cudaGetLastError());
}

void LMFitCUDA::calc_curve_values()
{
    dim3  threads(1, 1, 1);
    dim3  blocks(1, 1, 1);

    threads.x = info_.n_points_ * info_.n_fits_per_block_ / info_.n_blocks_per_fit_;

    if (info_.n_blocks_per_fit_ > 1)
        threads.x += info_.n_points_ % threads.x;

    blocks.x = n_fits_ / info_.n_fits_per_block_ * info_.n_blocks_per_fit_;

    cuda_calc_curve_values <<< blocks, threads >>>(
        gpu_data_.parameters_,
        n_fits_,
        info_.n_points_,
        info_.n_parameters_,
        gpu_data_.finished_,
        gpu_data_.values_,
        gpu_data_.derivatives_,
        info_.n_fits_per_block_,
        info_.n_blocks_per_fit_,
        info_.model_id_,
        gpu_data_.chunk_index_,
        gpu_data_.user_info_,
        info_.user_info_size_);
    CUDA_CHECK_STATUS(cudaGetLastError());
}

void LMFitCUDA::calc_chi_squares()
{
    dim3  threads(1, 1, 1);
    dim3  blocks(1, 1, 1);

    threads.x = info_.power_of_two_n_points_ * info_.n_fits_per_block_ / info_.n_blocks_per_fit_;
    blocks.x = n_fits_ / info_.n_fits_per_block_ * info_.n_blocks_per_fit_;

    int const shared_size = sizeof(float) * threads.x;

    float * chi_squares = 
        info_.n_blocks_per_fit_ > 1 ? gpu_data_.subtotals_ : gpu_data_.chi_squares_;

    cuda_calculate_chi_squares <<< blocks, threads, shared_size >>>(
        chi_squares,
        gpu_data_.states_,
        gpu_data_.data_,
        gpu_data_.values_,
        gpu_data_.weights_,
        info_.n_points_,
        n_fits_,
        info_.estimator_id_,
        gpu_data_.finished_,
        info_.n_fits_per_block_,
{'fraction_non_alphanumeric': 0.09536518103371586, 'fraction_numerical': 0.013699848767903212, 'mean_word_length': 3.0380747126436782, 'pattern_counts': {'":': 0, '<': 43, '<?xml version=': 0, '>': 45, 'https://': 0, 'lorem ipsum': 0, 'www.': 0, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 0, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '2842094', 'n_tokens_mistral': 4318, 'n_tokens_neox': 3970, 'n_words': 662}