From 705bd0c0841b3c3b6a24010f9ae2f6d7ce6346e8 Mon Sep 17 00:00:00 2001 From: Kane Scipioni Date: Tue, 25 Feb 2025 06:30:11 -0600 Subject: [PATCH 01/39] added runtime cuda toggle --- dlib/cuda/tensor_tools.cpp | 754 +++++++++++++++++++++---------------- dlib/cuda/tensor_tools.h | 155 ++++++-- dlib/test/dnn.cpp | 2 + 3 files changed, 570 insertions(+), 341 deletions(-) diff --git a/dlib/cuda/tensor_tools.cpp b/dlib/cuda/tensor_tools.cpp index 8dece8369f..c62b110bdc 100644 --- a/dlib/cuda/tensor_tools.cpp +++ b/dlib/cuda/tensor_tools.cpp @@ -17,6 +17,17 @@ namespace dlib static std::atomic var(true); return var; } + + bool& use_cuda_impl ( + ) + { +#ifdef DLIB_USE_CUDA + thread_local bool var(true); +#else + thread_local bool var(false); +#endif + return var; + } } bool dnn_prefer_fastest_algorithms ( @@ -36,6 +47,19 @@ namespace dlib { dnn_prefer_fastest_algo() = false; } + + bool use_cuda( + ) + { + return use_cuda_impl(); + } + + void set_use_cuda( + bool flag + ) + { + use_cuda_impl() = flag; + } } namespace dlib { namespace tt @@ -50,10 +74,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::inverse_norms(invnorms, data, eps); -#else - invnorms = reciprocal(sqrt(sum_cols(squared(mat(data))) + eps)); + if (use_cuda()) + cuda::inverse_norms(invnorms, data, eps); + else #endif + invnorms = reciprocal(sqrt(sum_cols(squared(mat(data))) + eps)); } void dot_prods ( @@ -63,10 +88,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::dot_prods(out, lhs, rhs); -#else - out = sum_cols(pointwise_multiply(mat(lhs), mat(rhs))); + if (use_cuda()) + cuda::dot_prods(out, lhs, rhs); + else #endif + out = sum_cols(pointwise_multiply(mat(lhs), mat(rhs))); } void dot_prods ( @@ -77,12 +103,19 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::dot_prods(add_to, out, lhs, rhs); -#else - if (add_to) - out += sum_cols(pointwise_multiply(mat(lhs), mat(rhs))); + if (use_cuda()) + { + cuda::dot_prods(add_to, out, lhs, rhs); + } else - out = sum_cols(pointwise_multiply(mat(lhs), mat(rhs))); + { +#endif + if (add_to) + out += sum_cols(pointwise_multiply(mat(lhs), mat(rhs))); + else + out = sum_cols(pointwise_multiply(mat(lhs), mat(rhs))); +#ifdef DLIB_USE_CUDA + } #endif } @@ -100,10 +133,11 @@ namespace dlib { namespace tt DLIB_CASSERT(m.size()/m.num_samples() == v.size()); #ifdef DLIB_USE_CUDA - cuda::scale_columns(out, m, v); -#else - out = scale_columns(mat(m), mat(v)); + if (use_cuda()) + cuda::scale_columns(out, m, v); + else #endif + out = scale_columns(mat(m), mat(v)); } void scale_rows ( @@ -120,10 +154,11 @@ namespace dlib { namespace tt DLIB_CASSERT(m.num_samples() == static_cast(v.size())); #ifdef DLIB_USE_CUDA - cuda::scale_rows(out, m, v); -#else - out = scale_rows(mat(m), mat(v)); + if (use_cuda()) + cuda::scale_rows(out, m, v); + else #endif + out = scale_rows(mat(m), mat(v)); } void scale_rows2 ( @@ -142,12 +177,19 @@ namespace dlib { namespace tt DLIB_CASSERT(static_cast(v1.size()) == m1.num_samples()); #ifdef DLIB_USE_CUDA - cuda::scale_rows2(beta, out, m1, m2, v1, v2); -#else - if (beta == 0) - out = scale_rows(mat(m1) - scale_rows(mat(m2),mat(v1)), mat(v2)); + if (use_cuda()) + { + cuda::scale_rows2(beta, out, m1, m2, v1, v2); + } else - out = beta*mat(out) + scale_rows(mat(m1) - scale_rows(mat(m2),mat(v1)), mat(v2)); + { +#endif + if (beta == 0) + out = scale_rows(mat(m1) - scale_rows(mat(m2),mat(v1)), mat(v2)); + else + out = beta*mat(out) + scale_rows(mat(m1) - scale_rows(mat(m2),mat(v1)), mat(v2)); +#ifdef DLIB_USE_CUDA + } #endif } @@ -161,10 +203,11 @@ namespace dlib { namespace tt DLIB_CASSERT(dest.size() == src.size()); #ifdef DLIB_USE_CUDA - cuda::exp(dest,src); -#else - dest = exp(mat(src)); + if (use_cuda()) + cuda::exp(dest,src); + else #endif + dest = exp(mat(src)); } // ---------------------------------------------------------------------------------------- @@ -177,10 +220,11 @@ namespace dlib { namespace tt DLIB_CASSERT(dest.size() == src.size()); #ifdef DLIB_USE_CUDA - cuda::log(dest,src); -#else - dest = log(mat(src)); + if (use_cuda()) + cuda::log(dest,src); + else #endif + dest = log(mat(src)); } // ---------------------------------------------------------------------------------------- @@ -193,10 +237,11 @@ namespace dlib { namespace tt DLIB_CASSERT(dest.size() == src.size()); #ifdef DLIB_USE_CUDA - cuda::log10(dest,src); -#else - dest = log10(mat(src)); + if (use_cuda()) + cuda::log10(dest,src); + else #endif + dest = log10(mat(src)); } // ---------------------------------------------------------------------------------------- @@ -213,94 +258,101 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::gemm(beta, dest, alpha, lhs, trans_lhs, rhs, trans_rhs, mode); -#else - if (mode == operation_mode::CHANNEL_WISE) + if (use_cuda()) + { + cuda::gemm(beta, dest, alpha, lhs, trans_lhs, rhs, trans_rhs, mode); + } + else { - if (beta != 0) +#endif + if (mode == operation_mode::CHANNEL_WISE) { - if (trans_lhs && trans_rhs) - dest = alpha * trans(mat(lhs)) * trans(mat(rhs)) + beta * mat(dest); - else if (!trans_lhs && trans_rhs) - dest = alpha * mat(lhs) * trans(mat(rhs)) + beta * mat(dest); - else if (trans_lhs && !trans_rhs) - dest = alpha * trans(mat(lhs)) * mat(rhs) + beta * mat(dest); + if (beta != 0) + { + if (trans_lhs && trans_rhs) + dest = alpha * trans(mat(lhs)) * trans(mat(rhs)) + beta * mat(dest); + else if (!trans_lhs && trans_rhs) + dest = alpha * mat(lhs) * trans(mat(rhs)) + beta * mat(dest); + else if (trans_lhs && !trans_rhs) + dest = alpha * trans(mat(lhs)) * mat(rhs) + beta * mat(dest); + else + dest = alpha * mat(lhs) * mat(rhs) + beta * mat(dest); + } else - dest = alpha * mat(lhs) * mat(rhs) + beta * mat(dest); + { + if (trans_lhs && trans_rhs) + dest = alpha * trans(mat(lhs)) * trans(mat(rhs)); + else if (!trans_lhs && trans_rhs) + dest = alpha * mat(lhs) * trans(mat(rhs)); + else if (trans_lhs && !trans_rhs) + dest = alpha * trans(mat(lhs)) * mat(rhs); + else + dest = alpha * mat(lhs) * mat(rhs); + } } - else + else if (mode == operation_mode::PLANE_WISE) { - if (trans_lhs && trans_rhs) - dest = alpha * trans(mat(lhs)) * trans(mat(rhs)); - else if (!trans_lhs && trans_rhs) - dest = alpha * mat(lhs) * trans(mat(rhs)); - else if (trans_lhs && !trans_rhs) - dest = alpha * trans(mat(lhs)) * mat(rhs); - else - dest = alpha * mat(lhs) * mat(rhs); - } - } - else if (mode == operation_mode::PLANE_WISE) - { - auto is_matrix = [](const auto& tensor) { - return ((tensor.num_samples() * tensor.k() == 1 && tensor.nr() * tensor.nc() > 1) || - (tensor.num_samples() * tensor.k() > 1 && tensor.nr() * tensor.nc() == 1)); - }; + auto is_matrix = [](const auto& tensor) { + return ((tensor.num_samples() * tensor.k() == 1 && tensor.nr() * tensor.nc() > 1) || + (tensor.num_samples() * tensor.k() > 1 && tensor.nr() * tensor.nc() == 1)); + }; - long num_samples = std::min({ lhs.num_samples(), rhs.num_samples(), dest.num_samples() }); - long num_channels = std::min({ lhs.k(), rhs.k(), dest.k() }); - const bool lhs_is_matrix = is_matrix(lhs), rhs_is_matrix = is_matrix(rhs), dest_is_matrix = is_matrix(dest); + long num_samples = std::min({ lhs.num_samples(), rhs.num_samples(), dest.num_samples() }); + long num_channels = std::min({ lhs.k(), rhs.k(), dest.k() }); + const bool lhs_is_matrix = is_matrix(lhs), rhs_is_matrix = is_matrix(rhs), dest_is_matrix = is_matrix(dest); - if (lhs_is_matrix && rhs_is_matrix && dest_is_matrix) { - num_samples = num_channels = 1; - } + if (lhs_is_matrix && rhs_is_matrix && dest_is_matrix) { + num_samples = num_channels = 1; + } - long lhs_rows = (lhs_is_matrix && lhs.num_samples() > 1) ? lhs.num_samples() : lhs.nr(); - long lhs_cols = (lhs_is_matrix && lhs.k() > 1) ? lhs.k() : lhs.nc(); - long rhs_rows = (rhs_is_matrix && rhs.num_samples() > 1) ? rhs.num_samples() : rhs.nr(); - long rhs_cols = (rhs_is_matrix && rhs.k() > 1) ? rhs.k() : rhs.nc(); - long dest_rows = (dest_is_matrix && dest.num_samples() > 1) ? dest.num_samples() : dest.nr(); - long dest_cols = (dest_is_matrix && dest.k() > 1) ? dest.k() : dest.nc(); + long lhs_rows = (lhs_is_matrix && lhs.num_samples() > 1) ? lhs.num_samples() : lhs.nr(); + long lhs_cols = (lhs_is_matrix && lhs.k() > 1) ? lhs.k() : lhs.nc(); + long rhs_rows = (rhs_is_matrix && rhs.num_samples() > 1) ? rhs.num_samples() : rhs.nr(); + long rhs_cols = (rhs_is_matrix && rhs.k() > 1) ? rhs.k() : rhs.nc(); + long dest_rows = (dest_is_matrix && dest.num_samples() > 1) ? dest.num_samples() : dest.nr(); + long dest_cols = (dest_is_matrix && dest.k() > 1) ? dest.k() : dest.nc(); - const size_t lhs_plane_size = lhs_rows * lhs_cols; - const size_t rhs_plane_size = rhs_rows * rhs_cols; - const size_t dest_plane_size = dest_rows * dest_cols; + const size_t lhs_plane_size = lhs_rows * lhs_cols; + const size_t rhs_plane_size = rhs_rows * rhs_cols; + const size_t dest_plane_size = dest_rows * dest_cols; - for (long b = 0; b < num_samples; ++b) - { - for (long c = 0; c < num_channels; ++c) + for (long b = 0; b < num_samples; ++b) { - auto lhs_slice = lhs_is_matrix ? alias_tensor(lhs_rows, lhs_cols)(lhs, 0) : - alias_tensor(lhs_rows, lhs_cols)(lhs, (b * num_channels + c) * lhs_plane_size); - auto rhs_slice = rhs_is_matrix ? alias_tensor(rhs_rows, rhs_cols)(rhs, 0) : - alias_tensor(rhs_rows, rhs_cols)(rhs, (b * num_channels + c) * rhs_plane_size); - auto dest_slice = dest_is_matrix ? alias_tensor(dest_rows, dest_cols)(dest, 0) : - alias_tensor(dest_rows, dest_cols)(dest, (b * num_channels + c) * dest_plane_size); - - if (beta != 0) + for (long c = 0; c < num_channels; ++c) { - if (trans_lhs && trans_rhs) - dest_slice = alpha * trans(mat(lhs_slice)) * trans(mat(rhs_slice)) + beta * mat(dest_slice); - else if (!trans_lhs && trans_rhs) - dest_slice = alpha * mat(lhs_slice) * trans(mat(rhs_slice)) + beta * mat(dest_slice); - else if (trans_lhs && !trans_rhs) - dest_slice = alpha * trans(mat(lhs_slice)) * mat(rhs_slice) + beta * mat(dest_slice); + auto lhs_slice = lhs_is_matrix ? alias_tensor(lhs_rows, lhs_cols)(lhs, 0) : + alias_tensor(lhs_rows, lhs_cols)(lhs, (b * num_channels + c) * lhs_plane_size); + auto rhs_slice = rhs_is_matrix ? alias_tensor(rhs_rows, rhs_cols)(rhs, 0) : + alias_tensor(rhs_rows, rhs_cols)(rhs, (b * num_channels + c) * rhs_plane_size); + auto dest_slice = dest_is_matrix ? alias_tensor(dest_rows, dest_cols)(dest, 0) : + alias_tensor(dest_rows, dest_cols)(dest, (b * num_channels + c) * dest_plane_size); + + if (beta != 0) + { + if (trans_lhs && trans_rhs) + dest_slice = alpha * trans(mat(lhs_slice)) * trans(mat(rhs_slice)) + beta * mat(dest_slice); + else if (!trans_lhs && trans_rhs) + dest_slice = alpha * mat(lhs_slice) * trans(mat(rhs_slice)) + beta * mat(dest_slice); + else if (trans_lhs && !trans_rhs) + dest_slice = alpha * trans(mat(lhs_slice)) * mat(rhs_slice) + beta * mat(dest_slice); + else + dest_slice = alpha * mat(lhs_slice) * mat(rhs_slice) + beta * mat(dest_slice); + } else - dest_slice = alpha * mat(lhs_slice) * mat(rhs_slice) + beta * mat(dest_slice); - } - else - { - if (trans_lhs && trans_rhs) - dest_slice = alpha * trans(mat(lhs_slice)) * trans(mat(rhs_slice)); - else if (!trans_lhs && trans_rhs) - dest_slice = alpha * mat(lhs_slice) * trans(mat(rhs_slice)); - else if (trans_lhs && !trans_rhs) - dest_slice = alpha * trans(mat(lhs_slice)) * mat(rhs_slice); - else - dest_slice = alpha * mat(lhs_slice) * mat(rhs_slice); + { + if (trans_lhs && trans_rhs) + dest_slice = alpha * trans(mat(lhs_slice)) * trans(mat(rhs_slice)); + else if (!trans_lhs && trans_rhs) + dest_slice = alpha * mat(lhs_slice) * trans(mat(rhs_slice)); + else if (trans_lhs && !trans_rhs) + dest_slice = alpha * trans(mat(lhs_slice)) * mat(rhs_slice); + else + dest_slice = alpha * mat(lhs_slice) * mat(rhs_slice); + } } } } +#ifdef DLIB_USE_CUDA } #endif } @@ -313,10 +365,9 @@ namespace dlib { namespace tt unsigned long long seed ) #ifdef DLIB_USE_CUDA - :rnd(seed){} -#else - {rnd.set_seed(cast_to_string(seed)); } + :cuda_impl(seed) #endif + {cpu_impl.set_seed(cast_to_string(seed)); } void tensor_rand:: fill_gaussian ( @@ -327,11 +378,12 @@ namespace dlib { namespace tt { DLIB_CASSERT(data.size()%2 == 0); #ifdef DLIB_USE_CUDA - rnd.fill_gaussian(data, mean, stddev); -#else - for (auto& x : data) - x = rnd.get_random_gaussian()*stddev + mean; + if (use_cuda()) + cuda_impl.fill_gaussian(data, mean, stddev); + else #endif + for (auto& x : data) + x = cpu_impl.get_random_gaussian()*stddev + mean; } void tensor_rand:: @@ -340,11 +392,12 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - rnd.fill_uniform(data); -#else - for (auto& x : data) - x = rnd.get_random_float(); + if (use_cuda()) + cuda_impl.fill_uniform(data); + else #endif + for (auto& x : data) + x = cpu_impl.get_random_float(); } // ---------------------------------------------------------------------------------------- @@ -365,10 +418,11 @@ namespace dlib { namespace tt (src1.num_samples()==1 || src1.num_samples()==MD) && (src2.num_samples()==1 || src2.num_samples()==MD) ); #ifdef DLIB_USE_CUDA - cuda::multiply(add_to, dest, src1, src2); -#else - cpu::multiply(add_to, dest, src1, src2); + if (use_cuda()) + cuda::multiply(add_to, dest, src1, src2); + else #endif + cpu::multiply(add_to, dest, src1, src2); } @@ -380,10 +434,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::scale_channels(add_to, dest, src, scales); -#else - cpu::scale_channels(add_to, dest, src, scales); + if(use_cuda()) + cuda::scale_channels(add_to, dest, src, scales); + else #endif + cpu::scale_channels(add_to, dest, src, scales); } void multiply_conv ( @@ -394,10 +449,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::multiply_conv(add_to, dest, src1, src2); -#else - cpu::multiply_conv(add_to, dest, src1, src2); + if (use_cuda()) + cuda::multiply_conv(add_to, dest, src1, src2); + else #endif + cpu::multiply_conv(add_to, dest, src1, src2); } void multiply_zero_padded ( @@ -408,10 +464,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::multiply_zero_padded(add_to, dest, src1, src2); -#else - cpu::multiply_zero_padded(add_to, dest, src1, src2); + if (use_cuda()) + cuda::multiply_zero_padded(add_to, dest, src1, src2); + else #endif + cpu::multiply_zero_padded(add_to, dest, src1, src2); } // ---------------------------------------------------------------------------------------- @@ -424,10 +481,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::affine_transform(dest,src,A,B); -#else - cpu::affine_transform(dest,src,A,B); + if (use_cuda()) + cuda::affine_transform(dest,src,A,B); + else #endif + cpu::affine_transform(dest,src,A,B); } void affine_transform( @@ -437,10 +495,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::affine_transform(dest,src,A); -#else - cpu::affine_transform(dest,src,A,0); + if (use_cuda()) + cuda::affine_transform(dest,src,A); + else #endif + cpu::affine_transform(dest,src,A,0); } void affine_transform( @@ -453,10 +512,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::affine_transform(dest,src1,src2,A,B,C); -#else - cpu::affine_transform(dest,src1,src2,A,B,C); + if (use_cuda()) + cuda::affine_transform(dest,src1,src2,A,B,C); + else #endif + cpu::affine_transform(dest,src1,src2,A,B,C); } void affine_transform( @@ -468,10 +528,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::affine_transform(dest,src1,src2,A,B); -#else - cpu::affine_transform(dest,src1,src2,A,B,0); + if (use_cuda()) + cuda::affine_transform(dest,src1,src2,A,B); + else #endif + cpu::affine_transform(dest,src1,src2,A,B,0); } void affine_transform( @@ -486,10 +547,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::affine_transform(dest,src1,src2,src3,A,B,C,D); -#else - cpu::affine_transform(dest,src1,src2,src3,A,B,C,D); + if (use_cuda()) + cuda::affine_transform(dest,src1,src2,src3,A,B,C,D); + else #endif + cpu::affine_transform(dest,src1,src2,src3,A,B,C,D); } void affine_transform_range( @@ -505,10 +567,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::affine_transform_range(begin, end, dest,src1,src2,src3,A,B,C); -#else - cpu::affine_transform_range(begin, end, dest,src1,src2,src3,A,B,C); + if (use_cuda()) + cuda::affine_transform_range(begin, end, dest,src1,src2,src3,A,B,C); + else #endif + cpu::affine_transform_range(begin, end, dest,src1,src2,src3,A,B,C); } void affine_transform( @@ -523,10 +586,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::affine_transform(rect, dest,src1,src2,src3,A,B,C); -#else - cpu::affine_transform(rect, dest,src1,src2,src3,A,B,C); + if (use_cuda()) + cuda::affine_transform(rect, dest,src1,src2,src3,A,B,C); + else #endif + cpu::affine_transform(rect, dest,src1,src2,src3,A,B,C); } void affine_transform( @@ -540,10 +604,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::affine_transform_range(0,dest.size(),dest,src1,src2,src3,A,B,C); -#else - cpu::affine_transform_range(0,dest.size(),dest,src1,src2,src3,A,B,C); + if (use_cuda()) + cuda::affine_transform_range(0,dest.size(),dest,src1,src2,src3,A,B,C); + else #endif + cpu::affine_transform_range(0,dest.size(),dest,src1,src2,src3,A,B,C); } // ---------------------------------------------------------------------------------------- @@ -556,10 +621,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::affine_transform(dest,src,A,B); -#else - cpu::affine_transform(dest,src,A,B); + if (use_cuda()) + cuda::affine_transform(dest,src,A,B); + else #endif + cpu::affine_transform(dest,src,A,B); } // ---------------------------------------------------------------------------------------- @@ -572,10 +638,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::affine_transform_conv(dest,src,A,B); -#else - cpu::affine_transform_conv(dest,src,A,B); + if (use_cuda()) + cuda::affine_transform_conv(dest,src,A,B); + else #endif + cpu::affine_transform_conv(dest,src,A,B); } // ---------------------------------------------------------------------------------------- @@ -596,12 +663,13 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::compute_adam_update(begin, end, s, m, v, t, learning_rate, weight_decay, momentum1, - momentum2, params, params_grad); -#else - cpu::compute_adam_update(begin, end, s, m, v, t, learning_rate, weight_decay, momentum1, - momentum2, params, params_grad); + if (use_cuda()) + cuda::compute_adam_update(begin, end, s, m, v, t, learning_rate, weight_decay, momentum1, + momentum2, params, params_grad); + else #endif + cpu::compute_adam_update(begin, end, s, m, v, t, learning_rate, weight_decay, momentum1, + momentum2, params, params_grad); } // ---------------------------------------------------------------------------------------- @@ -617,10 +685,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::batch_normalize_inference(eps,dest,src,gamma,beta,running_means,running_variances); -#else - cpu::batch_normalize_inference(eps,dest,src,gamma,beta,running_means,running_variances); + if (use_cuda()) + cuda::batch_normalize_inference(eps,dest,src,gamma,beta,running_means,running_variances); + else #endif + cpu::batch_normalize_inference(eps,dest,src,gamma,beta,running_means,running_variances); } void batch_normalize ( @@ -637,10 +706,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::batch_normalize(eps,dest,means,vars,averaging_factor,running_means,running_variances,src,gamma,beta); -#else - cpu::batch_normalize(eps,dest,means,vars,averaging_factor,running_means,running_variances,src,gamma,beta); + if (use_cuda()) + cuda::batch_normalize(eps,dest,means,vars,averaging_factor,running_means,running_variances,src,gamma,beta); + else #endif + cpu::batch_normalize(eps,dest,means,vars,averaging_factor,running_means,running_variances,src,gamma,beta); } void batch_normalize_gradient ( @@ -657,10 +727,11 @@ namespace dlib { namespace tt { #ifdef DLIB_USE_CUDA - cuda::batch_normalize_gradient(eps,gradient_input, means, invstds, src, gamma, src_grad, gamma_grad, beta_grad); -#else - cpu::batch_normalize_gradient(eps,gradient_input, means, invstds, src, gamma, src_grad, gamma_grad, beta_grad); + if (use_cuda()) + cuda::batch_normalize_gradient(eps,gradient_input, means, invstds, src, gamma, src_grad, gamma_grad, beta_grad); + else #endif + cpu::batch_normalize_gradient(eps,gradient_input, means, invstds, src, gamma, src_grad, gamma_grad, beta_grad); } // ---------------------------------------------------------------------------------------- @@ -676,10 +747,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::batch_normalize_conv_inference(eps,dest,src,gamma,beta,running_means,running_variances); -#else - cpu::batch_normalize_conv_inference(eps,dest,src,gamma,beta,running_means,running_variances); + if (use_cuda()) + cuda::batch_normalize_conv_inference(eps,dest,src,gamma,beta,running_means,running_variances); + else #endif + cpu::batch_normalize_conv_inference(eps,dest,src,gamma,beta,running_means,running_variances); } void batch_normalize_conv ( @@ -696,10 +768,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::batch_normalize_conv(eps,dest,means,vars,averaging_factor,running_means,running_variances,src,gamma,beta); -#else - cpu::batch_normalize_conv(eps,dest,means,vars,averaging_factor,running_means,running_variances,src,gamma,beta); + if (use_cuda()) + cuda::batch_normalize_conv(eps,dest,means,vars,averaging_factor,running_means,running_variances,src,gamma,beta); + else #endif + cpu::batch_normalize_conv(eps,dest,means,vars,averaging_factor,running_means,running_variances,src,gamma,beta); } void batch_normalize_conv_gradient ( @@ -716,10 +789,11 @@ namespace dlib { namespace tt { #ifdef DLIB_USE_CUDA - cuda::batch_normalize_conv_gradient(eps,gradient_input, means, invstds, src, gamma, src_grad, gamma_grad, beta_grad); -#else - cpu::batch_normalize_conv_gradient(eps,gradient_input, means, invstds, src, gamma, src_grad, gamma_grad, beta_grad); + if (use_cuda()) + cuda::batch_normalize_conv_gradient(eps,gradient_input, means, invstds, src, gamma, src_grad, gamma_grad, beta_grad); + else #endif + cpu::batch_normalize_conv_gradient(eps,gradient_input, means, invstds, src, gamma, src_grad, gamma_grad, beta_grad); } // ---------------------------------------------------------------------------------------- @@ -735,10 +809,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::layer_normalize(eps, dest, means, vars, src, gamma, beta); -#else - cpu::layer_normalize(eps, dest, means, vars, src, gamma, beta); + if (use_cuda()) + cuda::layer_normalize(eps, dest, means, vars, src, gamma, beta); + else #endif + cpu::layer_normalize(eps, dest, means, vars, src, gamma, beta); } void layer_normalize_gradient ( @@ -756,10 +831,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::layer_normalize_gradient(eps, gradient_input, means, invstds, src, gamma, src_grad, gamma_grad, beta_grad, dmeans, dvars); -#else - cpu::layer_normalize_gradient(eps, gradient_input, means, invstds, src, gamma, src_grad, gamma_grad, beta_grad, dmeans, dvars); + if (use_cuda()) + cuda::layer_normalize_gradient(eps, gradient_input, means, invstds, src, gamma, src_grad, gamma_grad, beta_grad, dmeans, dvars); + else #endif + cpu::layer_normalize_gradient(eps, gradient_input, means, invstds, src, gamma, src_grad, gamma_grad, beta_grad, dmeans, dvars); } // ---------------------------------------------------------------------------------------- @@ -773,10 +849,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::rms_normalize(eps, dest, scale, src, gamma); -#else - cpu::rms_normalize(eps, dest, scale, src, gamma); + if (use_cuda()) + cuda::rms_normalize(eps, dest, scale, src, gamma); + else #endif + cpu::rms_normalize(eps, dest, scale, src, gamma); } void rms_normalize_gradient( @@ -790,10 +867,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::rms_normalize_gradient(gradient_input, scale, src, gamma, src_grad, gamma_grad, dscale); -#else - cpu::rms_normalize_gradient(gradient_input, scale, src, gamma, src_grad, gamma_grad, dscale); + if (use_cuda()) + cuda::rms_normalize_gradient(gradient_input, scale, src, gamma, src_grad, gamma_grad, dscale); + else #endif + cpu::rms_normalize_gradient(gradient_input, scale, src, gamma, src_grad, gamma_grad, dscale); } // ---------------------------------------------------------------------------------------- @@ -804,10 +882,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::threshold(data,thresh); -#else - cpu::threshold(data,thresh); + if (use_cuda()) + cuda::threshold(data,thresh); + else #endif + cpu::threshold(data,thresh); } void dot ( @@ -818,10 +897,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::dot(a,b,result,idx); -#else - cpu::dot(a,b,result,idx); + if (use_cuda()) + cuda::dot(a,b,result,idx); + else #endif + cpu::dot(a,b,result,idx); } // ---------------------------------------------------------------------------------------- @@ -834,10 +914,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::add(beta,dest,alpha,src); -#else - cpu::add(beta,dest,alpha,src); + if (use_cuda()) + cuda::add(beta,dest,alpha,src); + else #endif + cpu::add(beta,dest,alpha,src); } // ---------------------------------------------------------------------------------------- @@ -849,10 +930,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::add(dest, src1, src2); -#else - cpu::add(dest, src1, src2); + if (use_cuda()) + cuda::add(dest, src1, src2); + else #endif + cpu::add(dest, src1, src2); } // ---------------------------------------------------------------------------------------- @@ -863,10 +945,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::assign_conv_bias_gradient(grad,gradient_input); -#else - cpu::assign_conv_bias_gradient(grad,gradient_input); + if (use_cuda()) + cuda::assign_conv_bias_gradient(grad,gradient_input); + else #endif + cpu::assign_conv_bias_gradient(grad,gradient_input); } // ---------------------------------------------------------------------------------------- @@ -877,10 +960,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::assign_bias_gradient(grad,gradient_input); -#else - cpu::assign_bias_gradient(grad,gradient_input); + if (use_cuda()) + cuda::assign_bias_gradient(grad,gradient_input); + else #endif + cpu::assign_bias_gradient(grad,gradient_input); } // ---------------------------------------------------------------------------------------- @@ -892,10 +976,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::softmax(dest, src, mode); -#else - cpu::softmax(dest, src, mode); + if (use_cuda()) + cuda::softmax(dest, src, mode); + else #endif + cpu::softmax(dest, src, mode); } void softmax_gradient( @@ -906,10 +991,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::softmax_gradient(grad, dest, gradient_input, mode); -#else - cpu::softmax_gradient(grad, dest, gradient_input, mode); + if (use_cuda()) + cuda::softmax_gradient(grad, dest, gradient_input, mode); + else #endif + cpu::softmax_gradient(grad, dest, gradient_input, mode); } // ---------------------------------------------------------------------------------------- @@ -920,10 +1006,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::softmax_all(dest,src); -#else - cpu::softmax_all(dest,src); + if (use_cuda()) + cuda::softmax_all(dest,src); + else #endif + cpu::softmax_all(dest,src); } void softmax_all_gradient ( @@ -933,10 +1020,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::softmax_all_gradient(grad, dest, gradient_input); -#else - cpu::softmax_all_gradient(grad, dest, gradient_input); + if (use_cuda()) + cuda::softmax_all_gradient(grad, dest, gradient_input); + else #endif + cpu::softmax_all_gradient(grad, dest, gradient_input); } // ---------------------------------------------------------------------------------------- @@ -947,10 +1035,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::sigmoid(dest,src); -#else - cpu::sigmoid(dest,src); + if (use_cuda()) + cuda::sigmoid(dest,src); + else #endif + cpu::sigmoid(dest,src); } void sigmoid_gradient ( @@ -960,10 +1049,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::sigmoid_gradient(grad, dest, gradient_input); -#else - cpu::sigmoid_gradient(grad, dest, gradient_input); + if (use_cuda()) + cuda::sigmoid_gradient(grad, dest, gradient_input); + else #endif + cpu::sigmoid_gradient(grad, dest, gradient_input); } // ---------------------------------------------------------------------------------------- @@ -974,10 +1064,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::mish(dest,src); -#else - cpu::mish(dest,src); + if (use_cuda()) + cuda::mish(dest,src); + else #endif + cpu::mish(dest,src); } void mish_gradient ( @@ -987,10 +1078,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::mish_gradient(grad, src, gradient_input); -#else - cpu::mish_gradient(grad, src, gradient_input); + if (use_cuda()) + cuda::mish_gradient(grad, src, gradient_input); + else #endif + cpu::mish_gradient(grad, src, gradient_input); } // ---------------------------------------------------------------------------------------- @@ -1001,10 +1093,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::relu(dest,src); -#else - cpu::relu(dest,src); + if (use_cuda()) + cuda::relu(dest,src); + else #endif + cpu::relu(dest,src); } void relu_gradient ( @@ -1014,10 +1107,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::relu_gradient(grad, dest, gradient_input); -#else - cpu::relu_gradient(grad, dest, gradient_input); + if (use_cuda()) + cuda::relu_gradient(grad, dest, gradient_input); + else #endif + cpu::relu_gradient(grad, dest, gradient_input); } // ---------------------------------------------------------------------------------------- @@ -1029,10 +1123,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::prelu(dest, src, param); -#else - cpu::prelu(dest, src, param); + if (use_cuda()) + cuda::prelu(dest, src, param); + else #endif + cpu::prelu(dest, src, param); } void prelu_gradient ( @@ -1044,10 +1139,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::prelu_gradient(grad, src, gradient_input, param, params_grad); -#else - cpu::prelu_gradient(grad, src, gradient_input, param, params_grad); + if (use_cuda()) + cuda::prelu_gradient(grad, src, gradient_input, param, params_grad); + else #endif + cpu::prelu_gradient(grad, src, gradient_input, param, params_grad); } // ---------------------------------------------------------------------------------------- @@ -1059,10 +1155,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::leaky_relu(dest, src, alpha); -#else - cpu::leaky_relu(dest, src, alpha); + if (use_cuda()) + cuda::leaky_relu(dest, src, alpha); + else #endif + cpu::leaky_relu(dest, src, alpha); } void leaky_relu_gradient ( @@ -1073,10 +1170,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::leaky_relu_gradient(grad, dest, gradient_input, alpha); -#else - cpu::leaky_relu_gradient(grad, dest, gradient_input, alpha); + if (use_cuda()) + cuda::leaky_relu_gradient(grad, dest, gradient_input, alpha); + else #endif + cpu::leaky_relu_gradient(grad, dest, gradient_input, alpha); } // ---------------------------------------------------------------------------------------- @@ -1087,10 +1185,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::tanh(dest,src); -#else - cpu::tanh(dest,src); + if (use_cuda()) + cuda::tanh(dest,src); + else #endif + cpu::tanh(dest,src); } void tanh_gradient ( @@ -1100,10 +1199,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::tanh_gradient(grad, dest, gradient_input); -#else - cpu::tanh_gradient(grad, dest, gradient_input); + if (use_cuda()) + cuda::tanh_gradient(grad, dest, gradient_input); + else #endif + cpu::tanh_gradient(grad, dest, gradient_input); } // ---------------------------------------------------------------------------------------- @@ -1115,10 +1215,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::clipped_relu(dest, src, ceiling); -#else - cpu::clipped_relu(dest, src, ceiling); + if (use_cuda()) + cuda::clipped_relu(dest, src, ceiling); + else #endif + cpu::clipped_relu(dest, src, ceiling); } void clipped_relu_gradient ( @@ -1129,10 +1230,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::clipped_relu_gradient(grad, dest, gradient_input, ceiling); -#else - cpu::clipped_relu_gradient(grad, dest, gradient_input, ceiling); + if (use_cuda()) + cuda::clipped_relu_gradient(grad, dest, gradient_input, ceiling); + else #endif + cpu::clipped_relu_gradient(grad, dest, gradient_input, ceiling); } // ---------------------------------------------------------------------------------------- @@ -1144,10 +1246,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::elu(dest, src, alpha); -#else - cpu::elu(dest, src, alpha); + if (use_cuda()) + cuda::elu(dest, src, alpha); + else #endif + cpu::elu(dest, src, alpha); } void elu_gradient ( @@ -1158,10 +1261,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::elu_gradient(grad, dest, gradient_input, alpha); -#else - cpu::elu_gradient(grad, dest, gradient_input, alpha); + if (use_cuda()) + cuda::elu_gradient(grad, dest, gradient_input, alpha); + else #endif + cpu::elu_gradient(grad, dest, gradient_input, alpha); } // ---------------------------------------------------------------------------------------- @@ -1172,10 +1276,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::gelu(dest,src); -#else - cpu::gelu(dest,src); + if (use_cuda()) + cuda::gelu(dest,src); + else #endif + cpu::gelu(dest,src); } void gelu_gradient ( @@ -1185,10 +1290,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::gelu_gradient(grad, src, gradient_input); -#else - cpu::gelu_gradient(grad, src, gradient_input); + if (use_cuda()) + cuda::gelu_gradient(grad, src, gradient_input); + else #endif + cpu::gelu_gradient(grad, src, gradient_input); } // ---------------------------------------------------------------------------------------- @@ -1201,10 +1307,11 @@ namespace dlib { namespace tt { DLIB_CASSERT(beta > 0); #ifdef DLIB_USE_CUDA - cuda::smelu(dest, src, beta); -#else - cpu::smelu(dest, src, beta); + if (use_cuda()) + cuda::smelu(dest, src, beta); + else #endif + cpu::smelu(dest, src, beta); } void smelu_gradient ( @@ -1216,10 +1323,11 @@ namespace dlib { namespace tt { DLIB_CASSERT(beta > 0); #ifdef DLIB_USE_CUDA - cuda::smelu_gradient(grad, dest, gradient_input, beta); -#else - cpu::smelu_gradient(grad, dest, gradient_input, beta); + if (use_cuda()) + cuda::smelu_gradient(grad, dest, gradient_input, beta); + else #endif + cpu::smelu_gradient(grad, dest, gradient_input, beta); } // ---------------------------------------------------------------------------------------- @@ -1230,10 +1338,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::silu(dest,src); -#else - cpu::silu(dest,src); + if (use_cuda()) + cuda::silu(dest,src); + else #endif + cpu::silu(dest,src); } void silu_gradient ( @@ -1243,10 +1352,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::silu_gradient(grad, src, gradient_input); -#else - cpu::silu_gradient(grad, src, gradient_input); + if (use_cuda()) + cuda::silu_gradient(grad, src, gradient_input); + else #endif + cpu::silu_gradient(grad, src, gradient_input); } // ---------------------------------------------------------------------------------------- @@ -1261,10 +1371,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::resize_bilinear(dest,dest_row_stride,dest_channel_stride, src,src_row_stride,src_channel_stride); -#else - cpu::resize_bilinear(dest,dest_row_stride,dest_channel_stride, src,src_row_stride,src_channel_stride); + if (use_cuda()) + cuda::resize_bilinear(dest,dest_row_stride,dest_channel_stride, src,src_row_stride,src_channel_stride); + else #endif + cpu::resize_bilinear(dest,dest_row_stride,dest_channel_stride, src,src_row_stride,src_channel_stride); } void resize_bilinear_gradient ( @@ -1277,10 +1388,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::resize_bilinear_gradient(grad,grad_row_stride,grad_channel_stride, gradient_input,gradient_input_row_stride,gradient_input_channel_stride); -#else - cpu::resize_bilinear_gradient(grad,grad_row_stride,grad_channel_stride, gradient_input,gradient_input_row_stride,gradient_input_channel_stride); + if (use_cuda()) + cuda::resize_bilinear_gradient(grad,grad_row_stride,grad_channel_stride, gradient_input,gradient_input_row_stride,gradient_input_channel_stride); + else #endif + cpu::resize_bilinear_gradient(grad,grad_row_stride,grad_channel_stride, gradient_input,gradient_input_row_stride,gradient_input_channel_stride); } // ------------------------------------------------------------------------------------ @@ -1294,10 +1406,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::reorg(add_to, dest, row_stride, col_stride, src); -#else - cpu::reorg(add_to, dest, row_stride, col_stride, src); + if (use_cuda()) + cuda::reorg(add_to, dest, row_stride, col_stride, src); + else #endif + cpu::reorg(add_to, dest, row_stride, col_stride, src); } void reorg_gradient ( @@ -1309,10 +1422,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::reorg_gradient(add_to, grad, row_stride, col_stride, gradient_input); -#else - cpu::reorg_gradient(add_to, grad, row_stride, col_stride, gradient_input); + if (use_cuda()) + cuda::reorg_gradient(add_to, grad, row_stride, col_stride, gradient_input); + else #endif + cpu::reorg_gradient(add_to, grad, row_stride, col_stride, gradient_input); } // ------------------------------------------------------------------------------------ @@ -1327,10 +1441,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::copy_tensor(add_to, dest, dest_k_offset, src, src_k_offset, count_k); -#else - cpu::copy_tensor(add_to, dest, dest_k_offset, src, src_k_offset, count_k); + if (use_cuda()) + cuda::copy_tensor(add_to, dest, dest_k_offset, src, src_k_offset, count_k); + else #endif + cpu::copy_tensor(add_to, dest, dest_k_offset, src, src_k_offset, count_k); } // ---------------------------------------------------------------------------------------- @@ -1345,10 +1460,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::copy_tensor(add_to, dest, dk, dnr, dnc , src, sk, snr, snc, k, nr, nc); -#else - cpu::copy_tensor(add_to, dest, dk, dnr, dnc, src, sk, snr, snc, k, nr, nc); + if (use_cuda()) + cuda::copy_tensor(add_to, dest, dk, dnr, dnc , src, sk, snr, snc, k, nr, nc); + else #endif + cpu::copy_tensor(add_to, dest, dk, dnr, dnc, src, sk, snr, snc, k, nr, nc); } // ---------------------------------------------------------------------------------------- @@ -1360,10 +1476,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - finv(m,out); -#else - out = dlib::inv(mat(m)); + if (use_cuda()) + finv(m,out); + else #endif + out = dlib::inv(mat(m)); } // ---------------------------------------------------------------------------------------- @@ -1375,10 +1492,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::transpose(add_to, dest, src); -#else - cpu::transpose(add_to, dest, src); + if (use_cuda()) + cuda::transpose(add_to, dest, src); + else #endif + cpu::transpose(add_to, dest, src); } // ---------------------------------------------------------------------------------------- @@ -1390,10 +1508,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::embeddings(dest, src, embs); -#else - cpu::embeddings(dest, src, embs); + if (use_cuda()) + cuda::embeddings(dest, src, embs); + else #endif + cpu::embeddings(dest, src, embs); } void embeddings_gradient( @@ -1406,10 +1525,11 @@ namespace dlib { namespace tt ) { #ifdef DLIB_USE_CUDA - cuda::embeddings_gradient(prev, gradient_input, grads, freqs, learning_rate, scale); -#else - cpu::embeddings_gradient(prev, gradient_input, grads, freqs, learning_rate, scale); + if (use_cuda()) + cuda::embeddings_gradient(prev, gradient_input, grads, freqs, learning_rate, scale); + else #endif + cpu::embeddings_gradient(prev, gradient_input, grads, freqs, learning_rate, scale); } // ---------------------------------------------------------------------------------------- diff --git a/dlib/cuda/tensor_tools.h b/dlib/cuda/tensor_tools.h index 18a5564f98..8e10e0d5e7 100644 --- a/dlib/cuda/tensor_tools.h +++ b/dlib/cuda/tensor_tools.h @@ -20,6 +20,9 @@ namespace dlib bool dnn_prefer_fastest_algorithms(); void set_dnn_prefer_fastest_algorithms(); void set_dnn_prefer_smallest_algorithms(); + + bool use_cuda(); + void set_use_cuda(bool flag); } namespace dlib { namespace tt @@ -292,10 +295,9 @@ namespace dlib { namespace tt !*/ #ifdef DLIB_USE_CUDA - cuda::curand_generator rnd; -#else - dlib::rand rnd; + cuda::curand_generator cuda_impl; #endif + dlib::rand cpu_impl; }; // ---------------------------------------------------------------------------------------- @@ -1074,14 +1076,30 @@ namespace dlib { namespace tt tensor_conv() {} void clear( - ) { impl.clear(); } + ) + { +#ifdef DLIB_USE_CUDA + if (use_cuda()) + cuda_impl.clear(); + else +#endif + cpu_impl.clear(); + } void operator() ( const bool add_to_output, tensor& output, const tensor& data, const tensor& filters - ) { impl(add_to_output,output,data,filters); } + ) + { +#ifdef DLIB_USE_CUDA + if (use_cuda()) + cuda_impl(add_to_output,output,data,filters); + else +#endif + cpu_impl(add_to_output,output,data,filters); + } /*! requires - setup() has been called. Specifically, setup() has been called like this: @@ -1107,7 +1125,15 @@ namespace dlib { namespace tt resizable_tensor& output, const tensor& data, const tensor& filters - ) { impl(add_to_output,output,data,filters); } + ) + { +#ifdef DLIB_USE_CUDA + if (use_cuda()) + cuda_impl(add_to_output,output,data,filters); + else +#endif + cpu_impl(add_to_output,output,data,filters); + } /*! requires - setup() has been called. Specifically, setup() has been called like this: @@ -1135,7 +1161,15 @@ namespace dlib { namespace tt const tensor& filters, const tensor& biases, bool use_relu - ) { impl(add_to_output,output,data,filters,biases,use_relu); } + ) + { +#ifdef DLIB_USE_CUDA + if (use_cuda()) + cuda_impl(add_to_output,output,data,filters,biases,use_relu); + else +#endif + cpu_impl(add_to_output,output,data,filters,biases,use_relu); + } /*! requires - setup() has been called. Specifically, setup() has been called like this: @@ -1167,7 +1201,15 @@ namespace dlib { namespace tt const tensor& filters, const tensor& biases, bool use_relu - ) { impl(add_to_output,output,data,filters,biases,use_relu); } + ) + { +#ifdef DLIB_USE_CUDA + if (use_cuda()) + cuda_impl(add_to_output,output,data,filters,biases,use_relu); + else +#endif + cpu_impl(add_to_output,output,data,filters,biases,use_relu); + } /*! requires - setup() has been called. Specifically, setup() has been called like this: @@ -1195,7 +1237,15 @@ namespace dlib { namespace tt const tensor& gradient_input, const tensor& filters, tensor& data_gradient - ) { impl.get_gradient_for_data(add_to_output,gradient_input,filters,data_gradient); } + ) + { +#ifdef DLIB_USE_CUDA + if (use_cuda()) + cuda_impl.get_gradient_for_data(add_to_output,gradient_input,filters,data_gradient); + else +#endif + cpu_impl.get_gradient_for_data(add_to_output,gradient_input,filters,data_gradient); + } /*! requires - One of the following must be true: @@ -1230,7 +1280,15 @@ namespace dlib { namespace tt const tensor& gradient_input, const tensor& data, tensor& filters_gradient - ) { impl.get_gradient_for_filters(add_to_output,gradient_input,data,filters_gradient); } + ) + { +#ifdef DLIB_USE_CUDA + if (use_cuda()) + cuda_impl.get_gradient_for_filters(add_to_output,gradient_input,data,filters_gradient); + else +#endif + cpu_impl.get_gradient_for_filters(add_to_output,gradient_input,data,filters_gradient); + } /*! requires - One of the following must be true: @@ -1268,7 +1326,15 @@ namespace dlib { namespace tt int stride_x, int padding_y, int padding_x - ) {impl.setup(data,filters,stride_y,stride_x,padding_y,padding_x); } + ) + { +#ifdef DLIB_USE_CUDA + if (use_cuda()) + cuda_impl.setup(data,filters,stride_y,stride_x,padding_y,padding_x); + else +#endif + cpu_impl.setup(data,filters,stride_y,stride_x,padding_y,padding_x); + } /*! requires - filters.k() == data.k() @@ -1292,11 +1358,9 @@ namespace dlib { namespace tt private: #ifdef DLIB_USE_CUDA - cuda::tensor_conv impl; -#else - cpu::tensor_conv impl; + cuda::tensor_conv cuda_impl; #endif - + cpu::tensor_conv cpu_impl; }; // ---------------------------------------------------------------------------------------- @@ -1317,7 +1381,13 @@ namespace dlib { namespace tt ) = default; void clear( - ) { impl.clear(); } + ) + { + if (use_cuda()) + cuda_impl.clear(); + else + cpu_impl.clear(); + } void setup_max_pooling( int window_height, @@ -1326,7 +1396,15 @@ namespace dlib { namespace tt int stride_x, int padding_y, int padding_x - ) { impl.setup_max_pooling(window_height, window_width, stride_y, stride_x, padding_y, padding_x); } + ) + { +#ifdef DLIB_USE_CUDA + if (use_cuda()) + cuda_impl.setup_max_pooling(window_height, window_width, stride_y, stride_x, padding_y, padding_x); + else +#endif + cpu_impl.setup_max_pooling(window_height, window_width, stride_y, stride_x, padding_y, padding_x); + } /*! requires - window_height > 0 @@ -1347,7 +1425,15 @@ namespace dlib { namespace tt int stride_x, int padding_y, int padding_x - ) { impl.setup_avg_pooling(window_height, window_width, stride_y, stride_x, padding_y, padding_x); } + ) + { +#ifdef DLIB_USE_CUDA + if (use_cuda()) + cuda_impl.setup_avg_pooling(window_height, window_width, stride_y, stride_x, padding_y, padding_x); + else +#endif + cpu_impl.setup_avg_pooling(window_height, window_width, stride_y, stride_x, padding_y, padding_x); + } /*! requires - window_height > 0 @@ -1362,12 +1448,26 @@ namespace dlib { namespace tt !*/ bool does_max_pooling( - ) const { return impl.does_max_pooling(); } + ) const + { +#ifdef DLIB_USE_CUDA + if (use_cuda()) + return cuda_impl.does_max_pooling(); + else +#endif + return cpu_impl.does_max_pooling(); + } void operator() ( resizable_tensor& dest, const tensor& src - ) { impl(dest, src); } + ) + { + if (use_cuda()) + cuda_impl(dest, src); + else + cpu_impl(dest, src); + } /*! requires - is_same_object(dest,src) == false @@ -1395,7 +1495,15 @@ namespace dlib { namespace tt const tensor& dest, const tensor& src, tensor& grad - ) { impl.get_gradient(gradient_input, dest, src, grad); } + ) + { +#ifdef DLIB_USE_CUDA + if (use_cuda()) + cuda_impl.get_gradient(gradient_input, dest, src, grad); + else +#endif + cpu_impl.get_gradient(gradient_input, dest, src, grad); + } /*! requires - have_same_dimensions(gradient_input,dest) == true @@ -1413,10 +1521,9 @@ namespace dlib { namespace tt private: #ifdef DLIB_USE_CUDA - cuda::pooling impl; -#else - cpu::pooling impl; + cuda::pooling cuda_impl; #endif + cpu::pooling cpu_impl; }; // ---------------------------------------------------------------------------------------- diff --git a/dlib/test/dnn.cpp b/dlib/test/dnn.cpp index 9316a0edc6..1da3bac777 100644 --- a/dlib/test/dnn.cpp +++ b/dlib/test/dnn.cpp @@ -5026,6 +5026,7 @@ void test_multm_prev() test_tagging(); #ifdef DLIB_USE_CUDA + set_use_cuda(true); test_affine_rect(); test_conv(); test_more_ops2(); @@ -5046,6 +5047,7 @@ void test_multm_prev() test_copy_tensor_add_to_gpu(); test_scale_channels(); #endif + set_use_cuda(false); test_tensor_resize_bilinear(2, 3, 6,6, 11, 11); test_tensor_resize_bilinear(2, 3, 6,6, 3, 4); test_tensor_resize_bilinear(2, 3, 5,6, 12, 21); From 4b8feabebc2e3899c56acd373ff46ed36777294b Mon Sep 17 00:00:00 2001 From: Kane Scipioni Date: Tue, 25 Feb 2025 06:37:43 -0600 Subject: [PATCH 02/39] only set if built with cuda --- dlib/cuda/tensor_tools.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/dlib/cuda/tensor_tools.cpp b/dlib/cuda/tensor_tools.cpp index c62b110bdc..fcffb26c00 100644 --- a/dlib/cuda/tensor_tools.cpp +++ b/dlib/cuda/tensor_tools.cpp @@ -58,7 +58,9 @@ namespace dlib bool flag ) { +#ifdef DLIB_USE_CUDA use_cuda_impl() = flag; +#endif } } From 3ea1be177c73e101d3598a55e5335b640c09eea4 Mon Sep 17 00:00:00 2001 From: Kane Scipioni Date: Tue, 25 Feb 2025 08:43:51 -0600 Subject: [PATCH 03/39] fixed missing macro condition --- dlib/cuda/tensor_tools.h | 2 ++ 1 file changed, 2 insertions(+) diff --git a/dlib/cuda/tensor_tools.h b/dlib/cuda/tensor_tools.h index 8e10e0d5e7..2d535a2aad 100644 --- a/dlib/cuda/tensor_tools.h +++ b/dlib/cuda/tensor_tools.h @@ -1383,9 +1383,11 @@ namespace dlib { namespace tt void clear( ) { +#ifdef DLIB_USE_CUDA if (use_cuda()) cuda_impl.clear(); else +#endif cpu_impl.clear(); } From 9ca48095a21fb0f6acfda142b350b8987ec0d509 Mon Sep 17 00:00:00 2001 From: Kane Scipioni Date: Tue, 25 Feb 2025 20:21:48 -0600 Subject: [PATCH 04/39] Fixed another missing directive --- dlib/cuda/tensor_tools.h | 2 ++ 1 file changed, 2 insertions(+) diff --git a/dlib/cuda/tensor_tools.h b/dlib/cuda/tensor_tools.h index 2d535a2aad..06a7fc01da 100644 --- a/dlib/cuda/tensor_tools.h +++ b/dlib/cuda/tensor_tools.h @@ -1465,9 +1465,11 @@ namespace dlib { namespace tt const tensor& src ) { +#ifdef DLIB_USE_CUDA if (use_cuda()) cuda_impl(dest, src); else +#endif cpu_impl(dest, src); } /*! From 49f8b0a3bad078147cd87d03f8b1f71c630ca149 Mon Sep 17 00:00:00 2001 From: tabudz <64760144+tabudz@users.noreply.github.com> Date: Thu, 27 Feb 2025 11:47:32 +0800 Subject: [PATCH 05/39] Check image size when reading targa file (#3058) Throw an error when image width or height is 0. Fixes mozilla/mozjpeg#140, closes #7. --- dlib/external/libjpeg/rdtarga.c | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/dlib/external/libjpeg/rdtarga.c b/dlib/external/libjpeg/rdtarga.c index c72ad73f6c..efcf4a41d4 100644 --- a/dlib/external/libjpeg/rdtarga.c +++ b/dlib/external/libjpeg/rdtarga.c @@ -365,7 +365,8 @@ start_input_tga (j_compress_ptr cinfo, cjpeg_source_ptr sinfo) width <= 0 || height <= 0 || source->pixel_size < 1 || source->pixel_size > 4 || (UCH(targaheader[16]) & 7) != 0 || /* bits/pixel must be multiple of 8 */ - interlace_type != 0) /* currently don't allow interlaced image */ + interlace_type != 0 || /* currently don't allow interlaced image */ + width == 0 || height == 0) /* image width/height must be non-zero */ ERREXIT(cinfo, JERR_TGA_BADPARMS); if (subtype > 8) { From 2a82cc36914764f3ad4cbffb355a589e4305b0e2 Mon Sep 17 00:00:00 2001 From: Wenbing Li <10278425+wenbingl@users.noreply.github.com> Date: Wed, 26 Feb 2025 19:48:15 -0800 Subject: [PATCH 06/39] change the default u32string char type to char32_t (#3059) --- dlib/serialize.h | 1 + dlib/unicode/unicode.h | 2 +- 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/dlib/serialize.h b/dlib/serialize.h index 17004ef027..ed8dbd8033 100644 --- a/dlib/serialize.h +++ b/dlib/serialize.h @@ -626,6 +626,7 @@ namespace dlib USE_DEFAULT_INT_SERIALIZATION_FOR(unsigned long) USE_DEFAULT_INT_SERIALIZATION_FOR(uint64) USE_DEFAULT_INT_SERIALIZATION_FOR(int64) + USE_DEFAULT_INT_SERIALIZATION_FOR(char32_t) USE_DEFAULT_BYTE_SERIALIZATION_FOR(char) USE_DEFAULT_BYTE_SERIALIZATION_FOR(signed char) diff --git a/dlib/unicode/unicode.h b/dlib/unicode/unicode.h index c71e052682..ce199b2548 100644 --- a/dlib/unicode/unicode.h +++ b/dlib/unicode/unicode.h @@ -16,7 +16,7 @@ namespace dlib // ---------------------------------------------------------------------------------------- - using unichar = uint32; + using unichar = char32_t; using ustring = std::basic_string; // ---------------------------------------------------------------------------------------- From 8bf1f5232476cb0dabd0f28dfd7dfa0f4179163a Mon Sep 17 00:00:00 2001 From: Davis King Date: Mon, 3 Mar 2025 08:25:09 -0500 Subject: [PATCH 07/39] Change version --- dlib/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dlib/CMakeLists.txt b/dlib/CMakeLists.txt index dd67154f02..b82a9fc82a 100644 --- a/dlib/CMakeLists.txt +++ b/dlib/CMakeLists.txt @@ -18,7 +18,7 @@ project(dlib) set(CPACK_PACKAGE_NAME "dlib") set(CPACK_PACKAGE_VERSION_MAJOR "19") set(CPACK_PACKAGE_VERSION_MINOR "24") -set(CPACK_PACKAGE_VERSION_PATCH "99") +set(CPACK_PACKAGE_VERSION_PATCH "8") set(VERSION ${CPACK_PACKAGE_VERSION_MAJOR}.${CPACK_PACKAGE_VERSION_MINOR}.${CPACK_PACKAGE_VERSION_PATCH}) # Only print these messages once, even if dlib is added multiple times via add_subdirectory() if (NOT TARGET dlib) From 1a38b7aa89794f31187365784952fa20da1ee018 Mon Sep 17 00:00:00 2001 From: Davis King Date: Mon, 3 Mar 2025 08:25:37 -0500 Subject: [PATCH 08/39] change version back --- dlib/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dlib/CMakeLists.txt b/dlib/CMakeLists.txt index b82a9fc82a..dd67154f02 100644 --- a/dlib/CMakeLists.txt +++ b/dlib/CMakeLists.txt @@ -18,7 +18,7 @@ project(dlib) set(CPACK_PACKAGE_NAME "dlib") set(CPACK_PACKAGE_VERSION_MAJOR "19") set(CPACK_PACKAGE_VERSION_MINOR "24") -set(CPACK_PACKAGE_VERSION_PATCH "8") +set(CPACK_PACKAGE_VERSION_PATCH "99") set(VERSION ${CPACK_PACKAGE_VERSION_MAJOR}.${CPACK_PACKAGE_VERSION_MINOR}.${CPACK_PACKAGE_VERSION_PATCH}) # Only print these messages once, even if dlib is added multiple times via add_subdirectory() if (NOT TARGET dlib) From 4a402baedd2c2ecebe70a1c66c85baf93e0a77e7 Mon Sep 17 00:00:00 2001 From: Richard Barnes Date: Mon, 3 Mar 2025 11:54:17 -0800 Subject: [PATCH 09/39] `throw()` -> `noexcept` `throw()` has been deprecated since C++11 and is removed in C++20. `noexcept` is accepted alternative. --- dlib/any/storage.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dlib/any/storage.h b/dlib/any/storage.h index 10a7b85c9b..df0a8ea92c 100644 --- a/dlib/any/storage.h +++ b/dlib/any/storage.h @@ -25,7 +25,7 @@ namespace dlib !*/ public: - virtual const char * what() const throw() + virtual const char * what() const noexcept { return "bad_any_cast"; } From 3b5f484dc64ecdd852a3e30ba44b6ab0d6ef5d03 Mon Sep 17 00:00:00 2001 From: tabudz <64760144+tabudz@users.noreply.github.com> Date: Tue, 4 Mar 2025 19:15:50 +0700 Subject: [PATCH 10/39] Fix a bug when getting a gzip header extra field with inflate(). (#3063) If the extra field was larger than the space the user provided with inflateGetHeader(), and if multiple calls of inflate() delivered the extra header data, then there could be a buffer overflow of the provided space. This commit assures that provided space is not exceeded. --- dlib/external/zlib/inflate.c | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/dlib/external/zlib/inflate.c b/dlib/external/zlib/inflate.c index ac333e8c2e..d758a1cc86 100644 --- a/dlib/external/zlib/inflate.c +++ b/dlib/external/zlib/inflate.c @@ -758,9 +758,10 @@ int flush; copy = state->length; if (copy > have) copy = have; if (copy) { + len = state->head->extra_len - state->length; if (state->head != Z_NULL && - state->head->extra != Z_NULL) { - len = state->head->extra_len - state->length; + state->head->extra != Z_NULL && + len < state->head->extra_max) { zmemcpy(state->head->extra + len, next, len + copy > state->head->extra_max ? state->head->extra_max - len : copy); From 3b37d5c1f03c6dcc35f3e624983a66b3e5a1f7fc Mon Sep 17 00:00:00 2001 From: Sandro Date: Sun, 16 Mar 2025 19:44:15 +0100 Subject: [PATCH 11/39] Drop namespace std (#3067) Solve build error when compiling with GCC 15. --- tools/python/src/sequence_segmenter.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/tools/python/src/sequence_segmenter.cpp b/tools/python/src/sequence_segmenter.cpp index 9fde1e771d..d6959911c2 100644 --- a/tools/python/src/sequence_segmenter.cpp +++ b/tools/python/src/sequence_segmenter.cpp @@ -1,13 +1,13 @@ // Copyright (C) 2013 Davis E. King (davis@dlib.net) // License: Boost Software License See LICENSE.txt for the full license. +#include #include "opaque_types.h" #include #include #include using namespace dlib; -using namespace std; namespace py = pybind11; typedef matrix dense_vect; @@ -277,9 +277,9 @@ struct segmenter_params }; -string segmenter_params__str__(const segmenter_params& p) +std::string segmenter_params__str__(const segmenter_params& p) { - ostringstream sout; + std::ostringstream sout; if (p.use_BIO_model) sout << "BIO,"; else @@ -307,9 +307,9 @@ string segmenter_params__str__(const segmenter_params& p) return trim(sout.str()); } -string segmenter_params__repr__(const segmenter_params& p) +std::string segmenter_params__repr__(const segmenter_params& p) { - ostringstream sout; + std::ostringstream sout; sout << "<"; sout << segmenter_params__str__(p); sout << ">"; From f14e77b981649d9aebc18c52e24bab2152c9577e Mon Sep 17 00:00:00 2001 From: Cydral <53169060+Cydral@users.noreply.github.com> Date: Sun, 23 Mar 2025 21:46:23 +0100 Subject: [PATCH 12/39] Add Byte Pair Encoding (BPE) class for subword tokenization (#3056) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * Add new BPE_Tokenizer class to Dlib - Implement BPE (Byte Pair Encoding) tokenization - Add training and encoding methods - Include unit tests * Update * Update * Last update: optimize BPE tokenizer encoding with parallel paragraph processing * Use of “in-memory” files to avoid leaving any traces on disk during the test. * Add one DLIB_TEST_MSG() test per encoded/decoded string * Add bpe_tokenizer_abstract.h for documentation and integration --- dlib/test/queue.cpp | 12 +- dlib/test/static_set.cpp | 2 +- dlib/test/tokenizer.cpp | 62 +++- dlib/tokenizer.h | 2 +- dlib/tokenizer/bpe_tokenizer.h | 453 ++++++++++++++++++++++++ dlib/tokenizer/bpe_tokenizer_abstract.h | 155 ++++++++ tools/htmlify/htmlify.cpp | 4 +- 7 files changed, 676 insertions(+), 14 deletions(-) create mode 100644 dlib/tokenizer/bpe_tokenizer.h create mode 100644 dlib/tokenizer/bpe_tokenizer_abstract.h diff --git a/dlib/test/queue.cpp b/dlib/test/queue.cpp index efcdf6054f..de52f825a5 100644 --- a/dlib/test/queue.cpp +++ b/dlib/test/queue.cpp @@ -408,17 +408,17 @@ namespace dlog << LINFO << "testing sort_1a_c"; - queue_sort_test::sort_1a_c> (); + queue_sort_test::sort_1a_c>(); dlog << LINFO << "testing sort_1a"; - queue_sort_test::sort_1a>(); + queue_sort_test::sort_1a>(); dlog << LINFO << "testing sort_1b"; - queue_sort_test::sort_1b> (); + queue_sort_test::sort_1b>(); dlog << LINFO << "testing sort_1b_c"; - queue_sort_test::sort_1b_c>(); + queue_sort_test::sort_1b_c>(); dlog << LINFO << "testing sort_1c"; - queue_sort_test::sort_1c> (); + queue_sort_test::sort_1c>(); dlog << LINFO << "testing sort_1c_c"; - queue_sort_test::sort_1c_c>(); + queue_sort_test::sort_1c_c>(); } } a; diff --git a/dlib/test/static_set.cpp b/dlib/test/static_set.cpp index 0ad864e4ad..a97791e10c 100644 --- a/dlib/test/static_set.cpp +++ b/dlib/test/static_set.cpp @@ -39,7 +39,7 @@ namespace srand(static_cast(time(0))); - typedef queue::kernel_2a_c queue_of_int; + typedef dlib::queue::kernel_2a_c queue_of_int; typedef dlib::set::kernel_1a_c set_of_int; queue_of_int q, qb, qc; diff --git a/dlib/test/tokenizer.cpp b/dlib/test/tokenizer.cpp index 95a95a7e1d..92def09372 100644 --- a/dlib/test/tokenizer.cpp +++ b/dlib/test/tokenizer.cpp @@ -1,9 +1,10 @@ -// Copyright (C) 2005 Davis E. King (davis@dlib.net) +// Copyright (C) 2005 Davis E. King (davis@dlib.net) // License: Boost Software License See LICENSE.txt for the full license. #include #include +#include #include #include "tester.h" @@ -350,9 +351,62 @@ namespace } + string postprocess_decoded_text(const string& decoded) { + string result = decoded; + result = regex_replace(result, std::regex(""), ""); + result = regex_replace(result, std::regex(""), "\n"); + if (!result.empty() && result.back() == '\n') result.pop_back(); + return result; + } + + template < + typename bpe_tok + > + void bpe_tokenizer_test( + ) + /*! + requires + - bpe_tok is an implementation of bpe_tokenizer.h + ensures + - runs tests on bpe_tok for compliance with the specs + !*/ + { + print_spinner(); + + bpe_tok test; + + std::string training_text = R"( + Byte Pair Encoding (BPE) is a subword tokenization algorithm widely used in Natural Language Processing (NLP). + It iteratively merges the most frequent pairs of bytes or characters to form a vocabulary of subword units. + This approach is particularly useful for handling out-of-vocabulary words and reducing the size of the vocabulary + while maintaining the ability to represent any text. BPE was introduced in the paper "Neural Machine Translation + of Rare Words with Subword Units" by Sennrich et al. in 2016. The algorithm is simple yet effective and has been + adopted in many state-of-the-art NLP models, including GPT and BERT. + )"; + test.train(training_text, 300, true); + std::ostringstream out_stream; + serialize(test, out_stream); + bpe_tok loaded_test; + std::istringstream in_stream(out_stream.str()); + deserialize(loaded_test, in_stream); + + std::vector test_strings = { + u8"This is a test of the tokenisation process...\nimplemented in the Dlib library!", // English + u8"Ceci est un test du processus de\ntokenisation implémenté dans\nla bibliothèque Dlib!", // French + u8"Dette er en test af tokeniseringsprocessen implementeret i Dlib-biblioteket!", // Danish + u8"这是对Dlib库中实现的标记化过程的测试！" // Chinese + }; + + for (const auto& text : test_strings) { + std::vector encoded = loaded_test.encode(text); + std::string decoded = postprocess_decoded_text(loaded_test.decode(encoded)); + + DLIB_TEST_MSG(text == decoded, "decoded: " << decoded); + } + } class tokenizer_tester : public tester { @@ -370,9 +424,9 @@ namespace tokenizer_kernel_test (); dlog << LINFO << "testing kernel_1a_c"; tokenizer_kernel_test(); + dlog << LINFO << "testing bpe_tokenizer"; + bpe_tokenizer_test(); } } a; -} - - +} \ No newline at end of file diff --git a/dlib/tokenizer.h b/dlib/tokenizer.h index 01b6fcf83a..c52343e37b 100644 --- a/dlib/tokenizer.h +++ b/dlib/tokenizer.h @@ -5,7 +5,7 @@ #include "tokenizer/tokenizer_kernel_1.h" #include "tokenizer/tokenizer_kernel_c.h" - +#include "tokenizer/bpe_tokenizer.h" namespace dlib { diff --git a/dlib/tokenizer/bpe_tokenizer.h b/dlib/tokenizer/bpe_tokenizer.h new file mode 100644 index 0000000000..2ebf032c16 --- /dev/null +++ b/dlib/tokenizer/bpe_tokenizer.h @@ -0,0 +1,453 @@ +#ifndef DLIB_BPE_TOKENIZER_H +#define DLIB_BPE_TOKENIZER_H + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "../base64.h" +#include "../serialize.h" + +namespace dlib +{ + + class bpe_tokenizer + { + public: + bpe_tokenizer() : vocab_size(BASE_VOCAB_SIZE) + { + // Initialize the base vocabulary with single bytes + for (int i = 0; i < BASE_VOCAB_SIZE; ++i) + vocab[i] = std::vector{ static_cast(i) }; + + // Initialize special tokens with sequential IDs + special_tokens = + { + {"", BASE_VOCAB_SIZE}, + {"", BASE_VOCAB_SIZE + 1}, + {"", BASE_VOCAB_SIZE + 2}, + {"", BASE_VOCAB_SIZE + 3}, + {"", BASE_VOCAB_SIZE + 4}, + {"", BASE_VOCAB_SIZE + 5}, + {"", BASE_VOCAB_SIZE + 7}, + {"", BASE_VOCAB_SIZE + 9}, + {"", BASE_VOCAB_SIZE + 10}, + {"", BASE_VOCAB_SIZE + 11}, + {"", BASE_VOCAB_SIZE + 12}, + {"", BASE_VOCAB_SIZE + 13}, + {"", BASE_VOCAB_SIZE + 14}, + {"", BASE_VOCAB_SIZE + 15}, + {"", BASE_VOCAB_SIZE + 16}, + {"", BASE_VOCAB_SIZE + 17}, + {"", BASE_VOCAB_SIZE + 18}, + {"", BASE_VOCAB_SIZE + 19}, + {"", BASE_VOCAB_SIZE + 20}, + {"", BASE_VOCAB_SIZE + 21}, + {"", BASE_VOCAB_SIZE + 22}, + {"", BASE_VOCAB_SIZE + 23}, + {"", BASE_VOCAB_SIZE + 24}, + {"", BASE_VOCAB_SIZE + 25}, + {"", BASE_VOCAB_SIZE + 26}, + {"", BASE_VOCAB_SIZE + 27} + }; + + // Initialize the vector of special token IDs + for (const auto& token : special_tokens) + special_token_map[token.second] = token.first; + } + + // Train the tokenizer on the given text + void train(const std::string& text, int vocab_size, bool verbose = false) + { + assert(vocab_size >= BASE_VOCAB_SIZE); + this->vocab_size = vocab_size; + int num_merges = vocab_size - BASE_VOCAB_SIZE; + + // Convert text to byte IDs + std::vector ids; + for (char c : text) ids.push_back(static_cast(c)); + + // Perform BPE merges + for (int i = 0; i < num_merges; ++i) { + auto stats = get_stats(ids); + if (stats.empty()) break; + + // Find the most frequent pair that does not exceed MAX_TOKEN_LENGTH + auto pair = get_most_frequent_pair(stats); + + // Check if the resulting token would exceed MAX_TOKEN_LENGTH + size_t new_token_length = vocab[pair.first].size() + vocab[pair.second].size(); + if (new_token_length > MAX_TOKEN_LENGTH) { + if (verbose) + { + std::cout << "\r" + << std::setw(100) << std::flush + << "\rskipping merge " << std::to_string(i + 1) << "/" << std::to_string(num_merges) << ": (" + << std::to_string(pair.first) << "," << std::to_string(pair.second) << ") -> new token length " + << std::to_string(new_token_length) << " exceeds limit of " << std::to_string(MAX_TOKEN_LENGTH) + << std::flush; + } + continue; // Skip this merge + } + + int idx = (BASE_VOCAB_SIZE + (int)special_tokens.size()) + i; + ids = merge(ids, pair, idx); + merges[pair] = idx; + vocab[idx].insert(vocab[idx].end(), vocab[pair.first].begin(), vocab[pair.first].end()); + vocab[idx].insert(vocab[idx].end(), vocab[pair.second].begin(), vocab[pair.second].end()); + + if (verbose) + { + std::cout << "\r" + << std::setw(100) << std::flush + << "\rmerge " << std::to_string(i + 1) << "/" << std::to_string(num_merges) << ": (" + << std::to_string(pair.first) << "," << std::to_string(pair.second) << ") -> " << std::to_string(idx) + << " (" << bytes_to_string(vocab[idx]) << ") had " + << std::to_string(stats[pair]) << " occurrences" + << std::flush; + } + } + std::cout << "\ntraining done\n"; + } + + // Encode the given text into subword tokens + std::vector encode(const std::string& text) + { + std::vector result_ids; + std::mutex result_mutex; + + // Split the text into paragraphs based on newline characters + std::vector paragraphs; + size_t start = 0, end = text.find('\n'); + while (end != std::string::npos) { + std::string paragraph = text.substr(start, end - start); + if (!paragraph.empty()) paragraphs.push_back(paragraph); + start = end + 1; + end = text.find('\n', start); + } + // Add the last paragraph (if any) and only if it's not empty + if (start < text.size()) { + std::string paragraph = text.substr(start); + if (!paragraph.empty()) paragraphs.push_back(paragraph); + } + + // Function to encode a single paragraph + auto encode_paragraph = [this](const std::string& paragraph) -> std::vector { + std::vector ids; + ids.reserve(paragraph.size()); + for (char c : paragraph) ids.push_back(static_cast(c)); + + auto stats = get_stats(ids); + std::priority_queue>> pq; + for (const auto& stat : stats) { + const std::pair& pair = stat.first; + if (merges.count(pair)) pq.push({ merges.at(pair), pair }); + } + + while (!pq.empty()) { + const auto& top_element = pq.top(); + const std::pair& pair = top_element.second; + pq.pop(); + + bool pair_found = false; + for (size_t i = 0; i < ids.size() - 1; ++i) { + if (ids[i] == pair.first && ids[i + 1] == pair.second) { + pair_found = true; + break; + } + } + if (!pair_found) continue; + + int idx = merges.at(pair); + ids = merge(ids, pair, idx); + + stats = get_stats(ids); + for (const auto& stat : stats) { + const std::pair& new_pair = stat.first; + if (merges.count(new_pair)) pq.push({ merges.at(new_pair), new_pair }); + } + } + + return ids; + }; + + // Special case: if there's only one paragraph, no need for threads + int sot_tok = get_special_token_id(""); + int eot_tok = get_special_token_id(""); + if (paragraphs.size() == 1) { + std::vector paragraph_ids = encode_paragraph(paragraphs[0]); + result_ids.push_back(sot_tok); + result_ids.insert(result_ids.end(), paragraph_ids.begin(), paragraph_ids.end()); + result_ids.push_back(eot_tok); + return result_ids; + } + + // Launch encoding tasks in parallel for multiple paragraphs + std::vector>> futures; + for (const auto& paragraph : paragraphs) + futures.push_back(std::async(std::launch::async, encode_paragraph, paragraph)); + + // Collect results in order + for (auto& future : futures) { + std::vector paragraph_ids = future.get(); + std::lock_guard lock(result_mutex); + result_ids.push_back(sot_tok); + result_ids.insert(result_ids.end(), paragraph_ids.begin(), paragraph_ids.end()); + result_ids.push_back(eot_tok); + } + return result_ids; + } + + // Decode a single token ID back into text + std::string decode(int id, bool display_special_tokens = true) + { + return decode(std::vector({ id }), display_special_tokens); + } + + // Decode a sequence of token IDs back into text + std::string decode(const std::vector& ids, bool display_special_tokens = true) + { + std::vector bytes; + int vocab_size = static_cast(get_vocab_size()); + for (int id : ids) + { + if (id < vocab_size) + { + // Check if the ID is a special token + auto it = special_token_map.find(id); + if (it != special_token_map.end()) + { + // It's a special token, get the corresponding string + if (display_special_tokens) bytes.insert(bytes.end(), it->second.begin(), it->second.end()); + } + else + { + // It's a regular token, get the bytes from the vocabulary + auto& token = vocab.at(id); + bytes.insert(bytes.end(), token.begin(), token.end()); + } + } + } + return std::string(bytes.begin(), bytes.end()); + } + + // Save the tokenizer model and vocabulary to file + friend void serialize(const bpe_tokenizer& tok, std::ostream& out) + { + dlib::serialize("bpe_tokenizer_", out); + + //--- + int nb_merges = tok.merges.size(); + dlib::serialize(nb_merges, out); + for (int idx = (BASE_VOCAB_SIZE + (int)tok.special_tokens.size()); + idx < (tok.vocab_size + (int)tok.special_tokens.size()); ++idx) + { + for (const auto& merge_pair : tok.merges) + { + if (merge_pair.second == idx) + { + dlib::serialize(merge_pair.first.first, out); + dlib::serialize(merge_pair.first.second, out); + break; + } + } + } + + //--- + int nb_vocab = (int)tok.vocab.size(); + dlib::serialize(nb_vocab, out); + for (const auto& v : tok.vocab) + { + std::string token_str = tok.bytes_to_string(v.second); + dlib::serialize(token_str, out); + dlib::serialize(v.first, out); + } + } + + // Load the tokenizer model and vocabulary from file + friend void deserialize(bpe_tokenizer& tok, std::istream& in) { + std::string version; + dlib::deserialize(version, in); + if (version != "bpe_tokenizer_") + throw dlib::serialization_error("Unexpected version '" + version + "' found while deserializing dlib::bpe_tokenizer_."); + + //--- + int idx = BASE_VOCAB_SIZE + (int)tok.special_tokens.size(), nb_merges, nb_vocab, a, b; + tok.merges.clear(); + dlib::deserialize(nb_merges, in); + for (int m = 0; m < nb_merges; m++) + { + dlib::deserialize(a, in); + dlib::deserialize(b, in); + tok.merges[{a, b}] = idx; + idx++; + } + + //--- + std::string token_str; + tok.vocab.clear(); + dlib::deserialize(nb_vocab, in); + for (int v = 0; v < nb_vocab; v++) + { + dlib::deserialize(token_str, in); + dlib::deserialize(idx, in); + tok.vocab[idx] = tok.string_to_bytes(token_str); + } + } + + // Get the ID of a special token + int get_special_token_id(const std::string& token) const + { + auto it = special_tokens.find(token); + if (it != special_tokens.end()) return it->second; + throw std::runtime_error("Special token not found: " + token); + } + + // Get the total vocabulary size + size_t get_vocab_size(void) const + { + return (vocab.size() + special_tokens.size()); + } + + private: + std::map special_tokens; + std::unordered_map special_token_map; + std::map, int> merges; + std::map> vocab; + int vocab_size; + + static const size_t MAX_TOKEN_LENGTH = 8; + static const int BASE_VOCAB_SIZE = 256; + + // Get frequency statistics of adjacent token pairs + struct pair_hash { + template + std::size_t operator()(const std::pair& p) const + { + auto hash1 = std::hash{}(p.first); + auto hash2 = std::hash{}(p.second); + return hash1 ^ (hash2 << 1); + } + }; + std::unordered_map, int, pair_hash> get_stats(const std::vector& ids) + { + std::unordered_map, int, pair_hash> global_stats; + std::mutex global_stats_mutex; + + auto worker = [&](size_t start, size_t end) { + std::unordered_map, int, pair_hash> local_stats; + for (size_t i = start; i < end - 1 && i + 1 < ids.size(); ++i) + local_stats[{ids[i], ids[i + 1]}]++; + + std::lock_guard lock(global_stats_mutex); + for (const auto& pair : local_stats) + global_stats[pair.first] += pair.second; + }; + + size_t num_threads = std::thread::hardware_concurrency(); + size_t segment_size = ids.size() / num_threads; + std::vector threads; + + for (size_t t = 0; t < num_threads; ++t) + { + size_t start = t * segment_size; + size_t end = (t == num_threads - 1) ? ids.size() : start + segment_size; + threads.emplace_back(worker, start, end); + } + + for (auto& thread : threads) thread.join(); + + return global_stats; + } + + // Finds the most frequent pair of tokens in the given statistics map that does not exceed the maximum token length + std::pair get_most_frequent_pair(const std::unordered_map, int, pair_hash>& stats) { + std::pair best_pair = { -1, -1 }; // Initialize the best pair to an invalid value + double max_score = 0; // Initialize the maximum score to 0 + + // Iterate over all pairs in the statistics map + for (const auto& stat : stats) { + const std::pair& pair = stat.first; // Extract the token pair + int count = stat.second; // Extract the frequency count + + // Check if the new token formed by merging the pair would exceed the maximum allowed length + size_t new_token_length = vocab[pair.first].size() + vocab[pair.second].size(); + if (new_token_length > MAX_TOKEN_LENGTH) continue; // Skip this pair if it exceeds the maximum token length + + // Calculate the score for this pair (frequency * length_penalty) + double score = (size_t)count * (new_token_length > (MAX_TOKEN_LENGTH / 2) ? 1.75 : 1.0); + + // Update the best pair if the current pair has a higher score + if (score > max_score) + { + best_pair = pair; + max_score = score; + } + } + + return best_pair; // Return the pair with the highest score + } + + // Merge the most frequent pair in the token sequence + std::vector merge(std::vector& ids, const std::pair& pair, int idx) { + std::vector new_ids; + new_ids.reserve(ids.size()); // Reserve space to avoid reallocations + + for (size_t i = 0; i < ids.size(); ++i) + { + if (i < ids.size() - 1 && ids[i] == pair.first && ids[i + 1] == pair.second) + { + new_ids.push_back(idx); // Replace the pair with the new token ID + i++; // Skip the next token + } + else new_ids.push_back(ids[i]); // Keep the current token + } + + return new_ids; + } + + // Decode/Encode a base64 string to/from a UTF-8 string + static std::string base64_decode(const std::string& base64_str) + { + dlib::base64 decoder; + std::istringstream sin(base64_str); + std::ostringstream sout; + decoder.decode(sin, sout); + return sout.str(); + } + static std::string base64_encode(const std::string& input) { + dlib::base64 encoder; + std::istringstream sin(input); + std::ostringstream sout; + encoder.encode(sin, sout); + return sout.str(); + } + + // Convert a sequence of bytes to a readable string + static std::string bytes_to_string(const std::vector& bytes) + { + std::string data(bytes.begin(), bytes.end()); + return base64_encode(data); + } + + // Convert a string representation of bytes back to bytes + static std::vector string_to_bytes(const std::string& str) + { + std::string decoded = base64_decode(str); + return std::vector(decoded.begin(), decoded.end()); + } + }; + +} + + +#endif // DLIB_BPE_TOKENIZER_H \ No newline at end of file diff --git a/dlib/tokenizer/bpe_tokenizer_abstract.h b/dlib/tokenizer/bpe_tokenizer_abstract.h new file mode 100644 index 0000000000..b3666922a4 --- /dev/null +++ b/dlib/tokenizer/bpe_tokenizer_abstract.h @@ -0,0 +1,155 @@ +// Copyright (C) 2025 Davis E. King (davis@dlib.net) +// License: Boost Software License See LICENSE.txt for the full license. +#undef DLIB_BPE_TOKENIZER_ABSTRACT_ +#ifdef DLIB_BPE_TOKENIZER_ABSTRACT_ + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace dlib +{ + + class bpe_tokenizer + { + /*! + CLASS bpe_tokenizer + A Byte Pair Encoding (BPE) tokenizer for text processing. + + This class implements a Byte Pair Encoding (BPE) tokenizer, which is a subword + tokenization algorithm commonly used in natural language processing (NLP). The + BPE algorithm iteratively merges the most frequent pairs of bytes or characters + to form a vocabulary of subword units. This approach is particularly useful for + handling out-of-vocabulary words and reducing the size of the vocabulary while + maintaining the ability to represent any text. + + The tokenizer supports special tokens, which can be used to mark specific elements + in the text (e.g., ``, ``, ``, etc.). These special tokens are + treated as atomic units during tokenization and are not subject to further splitting. + + The class provides methods for training the tokenizer on a given text corpus, encoding + text into subword tokens, and decoding tokens back into text. The tokenizer can be + serialized and deserialized to/from a file, allowing for easy storage and reuse. + + INITIAL VALUE + - The base vocabulary is initialized with single-byte tokens (0-255). + - Special tokens are pre-defined and assigned IDs starting from 256. + - The maximum token length is set to 8 bytes. + + WHAT THIS OBJECT REPRESENTS + This object represents a BPE tokenizer capable of encoding and decoding text + using a learned subword vocabulary. It is designed to handle UTF-8 encoded text + and supports multi-threaded processing for efficient tokenization. + + REFERENCES + - Sennrich, R., Haddow, B., & Birch, A. (2016). Neural Machine Translation of + Rare Words with Subword Units. In Proceedings of the 54th Annual Meeting of + the Association for Computational Linguistics (ACL 2016). + !*/ + public: + bpe_tokenizer(); + /*! + ensures + - Initializes the tokenizer with a base vocabulary of single-byte tokens (0-255). + - Pre-defines special tokens and assigns them unique IDs starting from 256. + !*/ + + void train( + const std::string& text, + int vocab_size, + bool verbose = false + ); + /*! + requires + - vocab_size >= 256 + ensures + - Trains the tokenizer on the provided text corpus. + - Iteratively merges the most frequent pairs of tokens to form a subword vocabulary + of size `vocab_size`. + - If `verbose` is true, progress information is printed to the standard output. + !*/ + + std::vector encode( + const std::string& text + ); + /*! + ensures + - Encodes the input text into a sequence of subword tokens. + - Special tokens are automatically added to mark the beginning and end of paragraphs. + - Returns a vector of token IDs representing the encoded text. + !*/ + + std::string decode( + const std::vector& ids, + bool display_special_tokens = true + ); + /*! + ensures + - Decodes a sequence of token IDs back into a human-readable string. + - If `display_special_tokens` is true, special tokens are included in the output. + - Returns the decoded text as a UTF-8 encoded string. + !*/ + + void serialize( + const bpe_tokenizer& tok, + std::ostream& out + ); + /*! + ensures + - Serializes the tokenizer's vocabulary and merge operations to the output stream. + - The serialized data can be used to reconstruct the tokenizer later. + !*/ + + void deserialize( + bpe_tokenizer& tok, + std::istream& in + ); + /*! + ensures + - Deserializes the tokenizer's vocabulary and merge operations from the input stream. + - Restores the tokenizer to the state it was in when serialized. + !*/ + + int get_special_token_id( + const std::string& token + ) const; + /*! + ensures + - Returns the ID of the specified special token. + - Throws an exception if the token is not found in the special tokens map. + !*/ + + size_t get_vocab_size() const; + /*! + ensures + - Returns the total size of the vocabulary, including base tokens and special tokens. + !*/ + + private: + // Private implementation details + std::map special_tokens; + std::unordered_map special_token_map; + std::map, int> merges; + std::map> vocab; + int vocab_size; + + static const size_t MAX_TOKEN_LENGTH = 8; + static const int BASE_VOCAB_SIZE = 256; + + // Helper functions + std::unordered_map, int, pair_hash> get_stats(const std::vector& ids); + std::pair get_most_frequent_pair(const std::unordered_map, int, pair_hash>& stats); + std::vector merge(std::vector& ids, const std::pair& pair, int idx); + std::string bytes_to_string(const std::vector& bytes); + std::vector string_to_bytes(const std::string& str); + }; + +} + +#endif // DLIB_BPE_TOKENIZER_ABSTRACT_ \ No newline at end of file diff --git a/tools/htmlify/htmlify.cpp b/tools/htmlify/htmlify.cpp index 3d281c2f00..c05b2e53a3 100644 --- a/tools/htmlify/htmlify.cpp +++ b/tools/htmlify/htmlify.cpp @@ -20,8 +20,8 @@ typedef cpp_pretty_printer::kernel_1a cprinter; typedef cpp_pretty_printer::kernel_2a bprinter; typedef dlib::map::kernel_1a map_string_to_string; typedef dlib::set::kernel_1a set_of_string; -typedef queue::kernel_1a queue_of_files; -typedef queue::kernel_1a queue_of_dirs; +typedef dlib::queue::kernel_1a queue_of_files; +typedef dlib::queue::kernel_1a queue_of_dirs; void print_manual ( ); From 207c3d7c56ff8b1f8a15407061dac01ff5826a93 Mon Sep 17 00:00:00 2001 From: Davis King Date: Sun, 23 Mar 2025 20:44:38 -0400 Subject: [PATCH 13/39] cleanup serialization code and add missing fields --- dlib/test/tokenizer.cpp | 4 +- dlib/tokenizer/bpe_tokenizer.h | 87 +++++++--------------------------- 2 files changed, 19 insertions(+), 72 deletions(-) diff --git a/dlib/test/tokenizer.cpp b/dlib/test/tokenizer.cpp index 92def09372..f4300666a2 100644 --- a/dlib/test/tokenizer.cpp +++ b/dlib/test/tokenizer.cpp @@ -384,7 +384,7 @@ namespace adopted in many state-of-the-art NLP models, including GPT and BERT. )"; - test.train(training_text, 300, true); + test.train(training_text, 300); std::ostringstream out_stream; serialize(test, out_stream); @@ -429,4 +429,4 @@ namespace } } a; -} \ No newline at end of file +} diff --git a/dlib/tokenizer/bpe_tokenizer.h b/dlib/tokenizer/bpe_tokenizer.h index 2ebf032c16..99bfcf23ee 100644 --- a/dlib/tokenizer/bpe_tokenizer.h +++ b/dlib/tokenizer/bpe_tokenizer.h @@ -1,3 +1,5 @@ +// Copyright (C) 2025 Davis E. King (davis@dlib.net) +// License: Boost Software License See LICENSE.txt for the full license. #ifndef DLIB_BPE_TOKENIZER_H #define DLIB_BPE_TOKENIZER_H @@ -14,6 +16,7 @@ #include "../base64.h" #include "../serialize.h" +#include "bpe_tokenizer_abstract.h" namespace dlib { @@ -113,10 +116,9 @@ namespace dlib << std::to_string(pair.first) << "," << std::to_string(pair.second) << ") -> " << std::to_string(idx) << " (" << bytes_to_string(vocab[idx]) << ") had " << std::to_string(stats[pair]) << " occurrences" - << std::flush; + << std::endl; } } - std::cout << "\ntraining done\n"; } // Encode the given text into subword tokens @@ -243,65 +245,25 @@ namespace dlib // Save the tokenizer model and vocabulary to file friend void serialize(const bpe_tokenizer& tok, std::ostream& out) { - dlib::serialize("bpe_tokenizer_", out); - - //--- - int nb_merges = tok.merges.size(); - dlib::serialize(nb_merges, out); - for (int idx = (BASE_VOCAB_SIZE + (int)tok.special_tokens.size()); - idx < (tok.vocab_size + (int)tok.special_tokens.size()); ++idx) - { - for (const auto& merge_pair : tok.merges) - { - if (merge_pair.second == idx) - { - dlib::serialize(merge_pair.first.first, out); - dlib::serialize(merge_pair.first.second, out); - break; - } - } - } - - //--- - int nb_vocab = (int)tok.vocab.size(); - dlib::serialize(nb_vocab, out); - for (const auto& v : tok.vocab) - { - std::string token_str = tok.bytes_to_string(v.second); - dlib::serialize(token_str, out); - dlib::serialize(v.first, out); - } + serialize("bpe_tokenizer2_", out); + serialize(tok.special_tokens, out); + serialize(tok.special_token_map, out); + serialize(tok.merges, out); + serialize(tok.vocab, out); + serialize(tok.vocab_size, out); } // Load the tokenizer model and vocabulary from file friend void deserialize(bpe_tokenizer& tok, std::istream& in) { std::string version; dlib::deserialize(version, in); - if (version != "bpe_tokenizer_") + if (version != "bpe_tokenizer2_") throw dlib::serialization_error("Unexpected version '" + version + "' found while deserializing dlib::bpe_tokenizer_."); - - //--- - int idx = BASE_VOCAB_SIZE + (int)tok.special_tokens.size(), nb_merges, nb_vocab, a, b; - tok.merges.clear(); - dlib::deserialize(nb_merges, in); - for (int m = 0; m < nb_merges; m++) - { - dlib::deserialize(a, in); - dlib::deserialize(b, in); - tok.merges[{a, b}] = idx; - idx++; - } - - //--- - std::string token_str; - tok.vocab.clear(); - dlib::deserialize(nb_vocab, in); - for (int v = 0; v < nb_vocab; v++) - { - dlib::deserialize(token_str, in); - dlib::deserialize(idx, in); - tok.vocab[idx] = tok.string_to_bytes(token_str); - } + deserialize(tok.special_tokens, in); + deserialize(tok.special_token_map, in); + deserialize(tok.merges, in); + deserialize(tok.vocab, in); + deserialize(tok.vocab_size, in); } // Get the ID of a special token @@ -415,15 +377,6 @@ namespace dlib return new_ids; } - // Decode/Encode a base64 string to/from a UTF-8 string - static std::string base64_decode(const std::string& base64_str) - { - dlib::base64 decoder; - std::istringstream sin(base64_str); - std::ostringstream sout; - decoder.decode(sin, sout); - return sout.str(); - } static std::string base64_encode(const std::string& input) { dlib::base64 encoder; std::istringstream sin(input); @@ -439,15 +392,9 @@ namespace dlib return base64_encode(data); } - // Convert a string representation of bytes back to bytes - static std::vector string_to_bytes(const std::string& str) - { - std::string decoded = base64_decode(str); - return std::vector(decoded.begin(), decoded.end()); - } }; } -#endif // DLIB_BPE_TOKENIZER_H \ No newline at end of file +#endif // DLIB_BPE_TOKENIZER_H From 9ca200f1ff82b2fdbc12790034ab640900a2126d Mon Sep 17 00:00:00 2001 From: Davis King Date: Sun, 23 Mar 2025 21:23:33 -0400 Subject: [PATCH 14/39] Some more cleanup --- dlib/tokenizer/bpe_tokenizer.h | 20 ++++--- dlib/tokenizer/bpe_tokenizer_abstract.h | 80 ++++++++++--------------- 2 files changed, 41 insertions(+), 59 deletions(-) diff --git a/dlib/tokenizer/bpe_tokenizer.h b/dlib/tokenizer/bpe_tokenizer.h index 99bfcf23ee..6547ae4a58 100644 --- a/dlib/tokenizer/bpe_tokenizer.h +++ b/dlib/tokenizer/bpe_tokenizer.h @@ -71,7 +71,7 @@ namespace dlib // Train the tokenizer on the given text void train(const std::string& text, int vocab_size, bool verbose = false) { - assert(vocab_size >= BASE_VOCAB_SIZE); + DLIB_CASSERT(vocab_size >= BASE_VOCAB_SIZE); this->vocab_size = vocab_size; int num_merges = vocab_size - BASE_VOCAB_SIZE; @@ -122,7 +122,7 @@ namespace dlib } // Encode the given text into subword tokens - std::vector encode(const std::string& text) + std::vector encode(const std::string& text) const { std::vector result_ids; std::mutex result_mutex; @@ -210,13 +210,13 @@ namespace dlib } // Decode a single token ID back into text - std::string decode(int id, bool display_special_tokens = true) + std::string decode(int id, bool display_special_tokens = true) const { return decode(std::vector({ id }), display_special_tokens); } // Decode a sequence of token IDs back into text - std::string decode(const std::vector& ids, bool display_special_tokens = true) + std::string decode(const std::vector& ids, bool display_special_tokens = true) const { std::vector bytes; int vocab_size = static_cast(get_vocab_size()); @@ -275,7 +275,7 @@ namespace dlib } // Get the total vocabulary size - size_t get_vocab_size(void) const + size_t get_vocab_size() const { return (vocab.size() + special_tokens.size()); } @@ -300,7 +300,7 @@ namespace dlib return hash1 ^ (hash2 << 1); } }; - std::unordered_map, int, pair_hash> get_stats(const std::vector& ids) + std::unordered_map, int, pair_hash> get_stats(const std::vector& ids) const { std::unordered_map, int, pair_hash> global_stats; std::mutex global_stats_mutex; @@ -332,7 +332,8 @@ namespace dlib } // Finds the most frequent pair of tokens in the given statistics map that does not exceed the maximum token length - std::pair get_most_frequent_pair(const std::unordered_map, int, pair_hash>& stats) { + std::pair get_most_frequent_pair(const std::unordered_map, int, pair_hash>& stats) const + { std::pair best_pair = { -1, -1 }; // Initialize the best pair to an invalid value double max_score = 0; // Initialize the maximum score to 0 @@ -342,7 +343,7 @@ namespace dlib int count = stat.second; // Extract the frequency count // Check if the new token formed by merging the pair would exceed the maximum allowed length - size_t new_token_length = vocab[pair.first].size() + vocab[pair.second].size(); + size_t new_token_length = vocab.at(pair.first).size() + vocab.at(pair.second).size(); if (new_token_length > MAX_TOKEN_LENGTH) continue; // Skip this pair if it exceeds the maximum token length // Calculate the score for this pair (frequency * length_penalty) @@ -360,7 +361,8 @@ namespace dlib } // Merge the most frequent pair in the token sequence - std::vector merge(std::vector& ids, const std::pair& pair, int idx) { + std::vector merge(std::vector& ids, const std::pair& pair, int idx) const + { std::vector new_ids; new_ids.reserve(ids.size()); // Reserve space to avoid reallocations diff --git a/dlib/tokenizer/bpe_tokenizer_abstract.h b/dlib/tokenizer/bpe_tokenizer_abstract.h index b3666922a4..1fa421d13a 100644 --- a/dlib/tokenizer/bpe_tokenizer_abstract.h +++ b/dlib/tokenizer/bpe_tokenizer_abstract.h @@ -19,9 +19,7 @@ namespace dlib class bpe_tokenizer { /*! - CLASS bpe_tokenizer - A Byte Pair Encoding (BPE) tokenizer for text processing. - + WHAT THIS OBJECT REPRESENTS This class implements a Byte Pair Encoding (BPE) tokenizer, which is a subword tokenization algorithm commonly used in natural language processing (NLP). The BPE algorithm iteratively merges the most frequent pairs of bytes or characters @@ -37,21 +35,17 @@ namespace dlib text into subword tokens, and decoding tokens back into text. The tokenizer can be serialized and deserialized to/from a file, allowing for easy storage and reuse. - INITIAL VALUE - - The base vocabulary is initialized with single-byte tokens (0-255). - - Special tokens are pre-defined and assigned IDs starting from 256. - - The maximum token length is set to 8 bytes. - - WHAT THIS OBJECT REPRESENTS - This object represents a BPE tokenizer capable of encoding and decoding text - using a learned subword vocabulary. It is designed to handle UTF-8 encoded text - and supports multi-threaded processing for efficient tokenization. - REFERENCES - Sennrich, R., Haddow, B., & Birch, A. (2016). Neural Machine Translation of Rare Words with Subword Units. In Proceedings of the 54th Annual Meeting of the Association for Computational Linguistics (ACL 2016). + + INITIAL VALUE + - The base vocabulary is initialized with single-byte tokens (0-255). + - Special tokens are pre-defined and assigned IDs starting from 256. + - The maximum token length is set to 8 bytes. !*/ + public: bpe_tokenizer(); /*! @@ -77,7 +71,7 @@ namespace dlib std::vector encode( const std::string& text - ); + ) const; /*! ensures - Encodes the input text into a sequence of subword tokens. @@ -88,7 +82,7 @@ namespace dlib std::string decode( const std::vector& ids, bool display_special_tokens = true - ); + ) const; /*! ensures - Decodes a sequence of token IDs back into a human-readable string. @@ -96,24 +90,11 @@ namespace dlib - Returns the decoded text as a UTF-8 encoded string. !*/ - void serialize( - const bpe_tokenizer& tok, - std::ostream& out - ); - /*! - ensures - - Serializes the tokenizer's vocabulary and merge operations to the output stream. - - The serialized data can be used to reconstruct the tokenizer later. - !*/ - - void deserialize( - bpe_tokenizer& tok, - std::istream& in - ); + std::string decode(int id, bool display_special_tokens = true) const + { return decode(std::vector({ id }), display_special_tokens); } /*! ensures - - Deserializes the tokenizer's vocabulary and merge operations from the input stream. - - Restores the tokenizer to the state it was in when serialized. + - decode a single token back into text. !*/ int get_special_token_id( @@ -130,26 +111,25 @@ namespace dlib ensures - Returns the total size of the vocabulary, including base tokens and special tokens. !*/ - - private: - // Private implementation details - std::map special_tokens; - std::unordered_map special_token_map; - std::map, int> merges; - std::map> vocab; - int vocab_size; - - static const size_t MAX_TOKEN_LENGTH = 8; - static const int BASE_VOCAB_SIZE = 256; - - // Helper functions - std::unordered_map, int, pair_hash> get_stats(const std::vector& ids); - std::pair get_most_frequent_pair(const std::unordered_map, int, pair_hash>& stats); - std::vector merge(std::vector& ids, const std::pair& pair, int idx); - std::string bytes_to_string(const std::vector& bytes); - std::vector string_to_bytes(const std::string& str); }; + void serialize( + const bpe_tokenizer& tok, + std::ostream& out + ); + /*! + ensures + - Saves the entire state of tok to out. + !*/ + + void deserialize( + bpe_tokenizer& tok, + std::istream& in + ); + /*! + ensures + - Restores the state of a bpe_tokenizer from a serialized state. + !*/ } -#endif // DLIB_BPE_TOKENIZER_ABSTRACT_ \ No newline at end of file +#endif // DLIB_BPE_TOKENIZER_ABSTRACT_ From 44fbbeb07ffbd93f780dc2190ef71f10c843064f Mon Sep 17 00:00:00 2001 From: Dobatymo Date: Sat, 5 Apr 2025 19:47:41 +0800 Subject: [PATCH 15/39] fix SyntaxWarning: invalid escape sequence '\(' (#3069) --- setup.py | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/setup.py b/setup.py index f6a44b59c8..9c0e57347f 100644 --- a/setup.py +++ b/setup.py @@ -231,9 +231,9 @@ def num_available_cpu_cores(ram_per_build_process_in_gb): def read_version_from_cmakelists(cmake_file): """Read version information """ - major = re.findall("set\(CPACK_PACKAGE_VERSION_MAJOR.*\"(.*)\"", open(cmake_file).read())[0] - minor = re.findall("set\(CPACK_PACKAGE_VERSION_MINOR.*\"(.*)\"", open(cmake_file).read())[0] - patch = re.findall("set\(CPACK_PACKAGE_VERSION_PATCH.*\"(.*)\"", open(cmake_file).read())[0] + major = re.findall("set\\(CPACK_PACKAGE_VERSION_MAJOR.*\"(.*)\"", open(cmake_file).read())[0] + minor = re.findall("set\\(CPACK_PACKAGE_VERSION_MINOR.*\"(.*)\"", open(cmake_file).read())[0] + patch = re.findall("set\\(CPACK_PACKAGE_VERSION_PATCH.*\"(.*)\"", open(cmake_file).read())[0] return major + '.' + minor + '.' + patch def read_entire_file(fname): From 47af03f27ff8eae18dc45e27e0064c5464ea34c2 Mon Sep 17 00:00:00 2001 From: Kane Scipioni Date: Tue, 6 May 2025 07:57:59 -0500 Subject: [PATCH 16/39] Defined macros to wrap use_cuda() branches --- dlib/cuda/tensor_tools.cpp | 733 ++++++++++++++++++++----------------- dlib/cuda/tensor_tools.h | 138 ++++--- 2 files changed, 481 insertions(+), 390 deletions(-) diff --git a/dlib/cuda/tensor_tools.cpp b/dlib/cuda/tensor_tools.cpp index fcffb26c00..b9818b3c73 100644 --- a/dlib/cuda/tensor_tools.cpp +++ b/dlib/cuda/tensor_tools.cpp @@ -75,12 +75,13 @@ namespace dlib { namespace tt const double eps ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::inverse_norms(invnorms, data, eps); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( invnorms = reciprocal(sqrt(sum_cols(squared(mat(data))) + eps)); + ) } void dot_prods ( @@ -89,12 +90,13 @@ namespace dlib { namespace tt const tensor& rhs ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::dot_prods(out, lhs, rhs); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( out = sum_cols(pointwise_multiply(mat(lhs), mat(rhs))); + ) } void dot_prods ( @@ -104,21 +106,16 @@ namespace dlib { namespace tt const tensor& rhs ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) - { + IF_DLIB_USE_CUDA( cuda::dot_prods(add_to, out, lhs, rhs); - } - else - { -#endif + ) + + IF_DLIB_NOT_USE_CUDA( if (add_to) out += sum_cols(pointwise_multiply(mat(lhs), mat(rhs))); else out = sum_cols(pointwise_multiply(mat(lhs), mat(rhs))); -#ifdef DLIB_USE_CUDA - } -#endif + ) } void scale_columns ( @@ -134,12 +131,13 @@ namespace dlib { namespace tt DLIB_CASSERT(m.size() != 0); DLIB_CASSERT(m.size()/m.num_samples() == v.size()); -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::scale_columns(out, m, v); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( out = scale_columns(mat(m), mat(v)); + ) } void scale_rows ( @@ -155,12 +153,13 @@ namespace dlib { namespace tt DLIB_CASSERT(m.size() != 0); DLIB_CASSERT(m.num_samples() == static_cast(v.size())); -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::scale_rows(out, m, v); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( out = scale_rows(mat(m), mat(v)); + ) } void scale_rows2 ( @@ -178,21 +177,16 @@ namespace dlib { namespace tt DLIB_CASSERT(is_vector(mat(v1))); DLIB_CASSERT(static_cast(v1.size()) == m1.num_samples()); -#ifdef DLIB_USE_CUDA - if (use_cuda()) - { + IF_DLIB_USE_CUDA( cuda::scale_rows2(beta, out, m1, m2, v1, v2); - } - else - { -#endif + ) + + IF_DLIB_NOT_USE_CUDA( if (beta == 0) out = scale_rows(mat(m1) - scale_rows(mat(m2),mat(v1)), mat(v2)); else out = beta*mat(out) + scale_rows(mat(m1) - scale_rows(mat(m2),mat(v1)), mat(v2)); -#ifdef DLIB_USE_CUDA - } -#endif + ) } // ---------------------------------------------------------------------------------------- @@ -204,12 +198,13 @@ namespace dlib { namespace tt { DLIB_CASSERT(dest.size() == src.size()); -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::exp(dest,src); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( dest = exp(mat(src)); + ) } // ---------------------------------------------------------------------------------------- @@ -221,12 +216,13 @@ namespace dlib { namespace tt { DLIB_CASSERT(dest.size() == src.size()); -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::log(dest,src); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( dest = log(mat(src)); + ) } // ---------------------------------------------------------------------------------------- @@ -238,12 +234,13 @@ namespace dlib { namespace tt { DLIB_CASSERT(dest.size() == src.size()); -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::log10(dest,src); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( dest = log10(mat(src)); + ) } // ---------------------------------------------------------------------------------------- @@ -259,14 +256,11 @@ namespace dlib { namespace tt operation_mode mode ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) - { + IF_DLIB_USE_CUDA( cuda::gemm(beta, dest, alpha, lhs, trans_lhs, rhs, trans_rhs, mode); - } - else - { -#endif + ) + + IF_DLIB_NOT_USE_CUDA( if (mode == operation_mode::CHANNEL_WISE) { if (beta != 0) @@ -354,9 +348,7 @@ namespace dlib { namespace tt } } } -#ifdef DLIB_USE_CUDA - } -#endif + ) } // ---------------------------------------------------------------------------------------- @@ -379,13 +371,15 @@ namespace dlib { namespace tt ) { DLIB_CASSERT(data.size()%2 == 0); -#ifdef DLIB_USE_CUDA - if (use_cuda()) + + IF_DLIB_USE_CUDA( cuda_impl.fill_gaussian(data, mean, stddev); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( for (auto& x : data) x = cpu_impl.get_random_gaussian()*stddev + mean; + ) } void tensor_rand:: @@ -393,13 +387,14 @@ namespace dlib { namespace tt tensor& data ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda_impl.fill_uniform(data); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( for (auto& x : data) x = cpu_impl.get_random_float(); + ) } // ---------------------------------------------------------------------------------------- @@ -419,12 +414,14 @@ namespace dlib { namespace tt DLIB_CASSERT((dest.num_samples()==1 || dest.num_samples()==MD) && (src1.num_samples()==1 || src1.num_samples()==MD) && (src2.num_samples()==1 || src2.num_samples()==MD) ); -#ifdef DLIB_USE_CUDA - if (use_cuda()) + + IF_DLIB_USE_CUDA( cuda::multiply(add_to, dest, src1, src2); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::multiply(add_to, dest, src1, src2); + ) } @@ -435,12 +432,13 @@ namespace dlib { namespace tt const tensor& scales ) { -#ifdef DLIB_USE_CUDA - if(use_cuda()) + IF_DLIB_USE_CUDA( cuda::scale_channels(add_to, dest, src, scales); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::scale_channels(add_to, dest, src, scales); + ) } void multiply_conv ( @@ -450,12 +448,13 @@ namespace dlib { namespace tt const tensor& src2 ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::multiply_conv(add_to, dest, src1, src2); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::multiply_conv(add_to, dest, src1, src2); + ) } void multiply_zero_padded ( @@ -465,12 +464,13 @@ namespace dlib { namespace tt const tensor& src2 ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::multiply_zero_padded(add_to, dest, src1, src2); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::multiply_zero_padded(add_to, dest, src1, src2); + ) } // ---------------------------------------------------------------------------------------- @@ -482,12 +482,13 @@ namespace dlib { namespace tt const float B ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::affine_transform(dest,src,A,B); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::affine_transform(dest,src,A,B); + ) } void affine_transform( @@ -496,12 +497,13 @@ namespace dlib { namespace tt const float A ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::affine_transform(dest,src,A); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::affine_transform(dest,src,A,0); + ) } void affine_transform( @@ -513,12 +515,13 @@ namespace dlib { namespace tt const float C ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::affine_transform(dest,src1,src2,A,B,C); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::affine_transform(dest,src1,src2,A,B,C); + ) } void affine_transform( @@ -529,12 +532,13 @@ namespace dlib { namespace tt const float B ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::affine_transform(dest,src1,src2,A,B); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::affine_transform(dest,src1,src2,A,B,0); + ) } void affine_transform( @@ -548,12 +552,13 @@ namespace dlib { namespace tt const float D ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::affine_transform(dest,src1,src2,src3,A,B,C,D); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::affine_transform(dest,src1,src2,src3,A,B,C,D); + ) } void affine_transform_range( @@ -568,12 +573,13 @@ namespace dlib { namespace tt const float C ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::affine_transform_range(begin, end, dest,src1,src2,src3,A,B,C); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::affine_transform_range(begin, end, dest,src1,src2,src3,A,B,C); + ) } void affine_transform( @@ -587,12 +593,13 @@ namespace dlib { namespace tt float C ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::affine_transform(rect, dest,src1,src2,src3,A,B,C); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::affine_transform(rect, dest,src1,src2,src3,A,B,C); + ) } void affine_transform( @@ -605,12 +612,13 @@ namespace dlib { namespace tt const float C ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::affine_transform_range(0,dest.size(),dest,src1,src2,src3,A,B,C); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::affine_transform_range(0,dest.size(),dest,src1,src2,src3,A,B,C); + ) } // ---------------------------------------------------------------------------------------- @@ -622,12 +630,13 @@ namespace dlib { namespace tt const tensor& B ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::affine_transform(dest,src,A,B); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::affine_transform(dest,src,A,B); + ) } // ---------------------------------------------------------------------------------------- @@ -639,12 +648,13 @@ namespace dlib { namespace tt const tensor& B ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::affine_transform_conv(dest,src,A,B); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::affine_transform_conv(dest,src,A,B); + ) } // ---------------------------------------------------------------------------------------- @@ -664,14 +674,15 @@ namespace dlib { namespace tt const tensor& params_grad ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::compute_adam_update(begin, end, s, m, v, t, learning_rate, weight_decay, momentum1, momentum2, params, params_grad); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::compute_adam_update(begin, end, s, m, v, t, learning_rate, weight_decay, momentum1, momentum2, params, params_grad); + ) } // ---------------------------------------------------------------------------------------- @@ -686,12 +697,13 @@ namespace dlib { namespace tt const tensor& running_variances ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::batch_normalize_inference(eps,dest,src,gamma,beta,running_means,running_variances); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::batch_normalize_inference(eps,dest,src,gamma,beta,running_means,running_variances); + ) } void batch_normalize ( @@ -707,12 +719,13 @@ namespace dlib { namespace tt const tensor& beta ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::batch_normalize(eps,dest,means,vars,averaging_factor,running_means,running_variances,src,gamma,beta); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::batch_normalize(eps,dest,means,vars,averaging_factor,running_means,running_variances,src,gamma,beta); + ) } void batch_normalize_gradient ( @@ -728,12 +741,13 @@ namespace dlib { namespace tt ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::batch_normalize_gradient(eps,gradient_input, means, invstds, src, gamma, src_grad, gamma_grad, beta_grad); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::batch_normalize_gradient(eps,gradient_input, means, invstds, src, gamma, src_grad, gamma_grad, beta_grad); + ) } // ---------------------------------------------------------------------------------------- @@ -748,12 +762,13 @@ namespace dlib { namespace tt const tensor& running_variances ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::batch_normalize_conv_inference(eps,dest,src,gamma,beta,running_means,running_variances); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::batch_normalize_conv_inference(eps,dest,src,gamma,beta,running_means,running_variances); + ) } void batch_normalize_conv ( @@ -769,12 +784,13 @@ namespace dlib { namespace tt const tensor& beta ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::batch_normalize_conv(eps,dest,means,vars,averaging_factor,running_means,running_variances,src,gamma,beta); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::batch_normalize_conv(eps,dest,means,vars,averaging_factor,running_means,running_variances,src,gamma,beta); + ) } void batch_normalize_conv_gradient ( @@ -790,12 +806,13 @@ namespace dlib { namespace tt ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::batch_normalize_conv_gradient(eps,gradient_input, means, invstds, src, gamma, src_grad, gamma_grad, beta_grad); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::batch_normalize_conv_gradient(eps,gradient_input, means, invstds, src, gamma, src_grad, gamma_grad, beta_grad); + ) } // ---------------------------------------------------------------------------------------- @@ -810,12 +827,13 @@ namespace dlib { namespace tt const tensor& beta ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::layer_normalize(eps, dest, means, vars, src, gamma, beta); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::layer_normalize(eps, dest, means, vars, src, gamma, beta); + ) } void layer_normalize_gradient ( @@ -832,12 +850,13 @@ namespace dlib { namespace tt resizable_tensor& dvars ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::layer_normalize_gradient(eps, gradient_input, means, invstds, src, gamma, src_grad, gamma_grad, beta_grad, dmeans, dvars); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::layer_normalize_gradient(eps, gradient_input, means, invstds, src, gamma, src_grad, gamma_grad, beta_grad, dmeans, dvars); + ) } // ---------------------------------------------------------------------------------------- @@ -850,12 +869,13 @@ namespace dlib { namespace tt const tensor& gamma ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::rms_normalize(eps, dest, scale, src, gamma); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::rms_normalize(eps, dest, scale, src, gamma); + ) } void rms_normalize_gradient( @@ -868,12 +888,13 @@ namespace dlib { namespace tt resizable_tensor& dscale ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::rms_normalize_gradient(gradient_input, scale, src, gamma, src_grad, gamma_grad, dscale); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::rms_normalize_gradient(gradient_input, scale, src, gamma, src_grad, gamma_grad, dscale); + ) } // ---------------------------------------------------------------------------------------- @@ -883,12 +904,13 @@ namespace dlib { namespace tt float thresh ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::threshold(data,thresh); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::threshold(data,thresh); + ) } void dot ( @@ -898,12 +920,13 @@ namespace dlib { namespace tt size_t idx ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::dot(a,b,result,idx); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::dot(a,b,result,idx); + ) } // ---------------------------------------------------------------------------------------- @@ -915,12 +938,13 @@ namespace dlib { namespace tt const tensor& src ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::add(beta,dest,alpha,src); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::add(beta,dest,alpha,src); + ) } // ---------------------------------------------------------------------------------------- @@ -931,12 +955,13 @@ namespace dlib { namespace tt const tensor& src2 ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::add(dest, src1, src2); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::add(dest, src1, src2); + ) } // ---------------------------------------------------------------------------------------- @@ -946,12 +971,13 @@ namespace dlib { namespace tt const tensor& gradient_input ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::assign_conv_bias_gradient(grad,gradient_input); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::assign_conv_bias_gradient(grad,gradient_input); + ) } // ---------------------------------------------------------------------------------------- @@ -961,12 +987,13 @@ namespace dlib { namespace tt const tensor& gradient_input ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::assign_bias_gradient(grad,gradient_input); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::assign_bias_gradient(grad,gradient_input); + ) } // ---------------------------------------------------------------------------------------- @@ -977,12 +1004,13 @@ namespace dlib { namespace tt operation_mode mode ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::softmax(dest, src, mode); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::softmax(dest, src, mode); + ) } void softmax_gradient( @@ -992,12 +1020,13 @@ namespace dlib { namespace tt operation_mode mode ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::softmax_gradient(grad, dest, gradient_input, mode); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::softmax_gradient(grad, dest, gradient_input, mode); + ) } // ---------------------------------------------------------------------------------------- @@ -1007,12 +1036,13 @@ namespace dlib { namespace tt const tensor& src ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::softmax_all(dest,src); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::softmax_all(dest,src); + ) } void softmax_all_gradient ( @@ -1021,12 +1051,13 @@ namespace dlib { namespace tt const tensor& gradient_input ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::softmax_all_gradient(grad, dest, gradient_input); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::softmax_all_gradient(grad, dest, gradient_input); + ) } // ---------------------------------------------------------------------------------------- @@ -1036,12 +1067,13 @@ namespace dlib { namespace tt const tensor& src ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::sigmoid(dest,src); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::sigmoid(dest,src); + ) } void sigmoid_gradient ( @@ -1050,12 +1082,13 @@ namespace dlib { namespace tt const tensor& gradient_input ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::sigmoid_gradient(grad, dest, gradient_input); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::sigmoid_gradient(grad, dest, gradient_input); + ) } // ---------------------------------------------------------------------------------------- @@ -1065,12 +1098,13 @@ namespace dlib { namespace tt const tensor& src ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::mish(dest,src); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::mish(dest,src); + ) } void mish_gradient ( @@ -1079,12 +1113,13 @@ namespace dlib { namespace tt const tensor& gradient_input ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::mish_gradient(grad, src, gradient_input); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::mish_gradient(grad, src, gradient_input); + ) } // ---------------------------------------------------------------------------------------- @@ -1094,12 +1129,13 @@ namespace dlib { namespace tt const tensor& src ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::relu(dest,src); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::relu(dest,src); + ) } void relu_gradient ( @@ -1108,12 +1144,13 @@ namespace dlib { namespace tt const tensor& gradient_input ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::relu_gradient(grad, dest, gradient_input); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::relu_gradient(grad, dest, gradient_input); + ) } // ---------------------------------------------------------------------------------------- @@ -1124,12 +1161,13 @@ namespace dlib { namespace tt const tensor& param ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::prelu(dest, src, param); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::prelu(dest, src, param); + ) } void prelu_gradient ( @@ -1140,12 +1178,13 @@ namespace dlib { namespace tt tensor& params_grad ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::prelu_gradient(grad, src, gradient_input, param, params_grad); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::prelu_gradient(grad, src, gradient_input, param, params_grad); + ) } // ---------------------------------------------------------------------------------------- @@ -1156,12 +1195,13 @@ namespace dlib { namespace tt const float alpha ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::leaky_relu(dest, src, alpha); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::leaky_relu(dest, src, alpha); + ) } void leaky_relu_gradient ( @@ -1171,12 +1211,13 @@ namespace dlib { namespace tt const float alpha ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::leaky_relu_gradient(grad, dest, gradient_input, alpha); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::leaky_relu_gradient(grad, dest, gradient_input, alpha); + ) } // ---------------------------------------------------------------------------------------- @@ -1186,12 +1227,13 @@ namespace dlib { namespace tt const tensor& src ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::tanh(dest,src); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::tanh(dest,src); + ) } void tanh_gradient ( @@ -1200,12 +1242,13 @@ namespace dlib { namespace tt const tensor& gradient_input ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::tanh_gradient(grad, dest, gradient_input); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::tanh_gradient(grad, dest, gradient_input); + ) } // ---------------------------------------------------------------------------------------- @@ -1216,12 +1259,13 @@ namespace dlib { namespace tt const float ceiling ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::clipped_relu(dest, src, ceiling); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::clipped_relu(dest, src, ceiling); + ) } void clipped_relu_gradient ( @@ -1231,12 +1275,13 @@ namespace dlib { namespace tt const float ceiling ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::clipped_relu_gradient(grad, dest, gradient_input, ceiling); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::clipped_relu_gradient(grad, dest, gradient_input, ceiling); + ) } // ---------------------------------------------------------------------------------------- @@ -1247,12 +1292,13 @@ namespace dlib { namespace tt const float alpha ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::elu(dest, src, alpha); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::elu(dest, src, alpha); + ) } void elu_gradient ( @@ -1262,12 +1308,13 @@ namespace dlib { namespace tt const float alpha ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::elu_gradient(grad, dest, gradient_input, alpha); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::elu_gradient(grad, dest, gradient_input, alpha); + ) } // ---------------------------------------------------------------------------------------- @@ -1277,12 +1324,13 @@ namespace dlib { namespace tt const tensor& src ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::gelu(dest,src); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::gelu(dest,src); + ) } void gelu_gradient ( @@ -1291,12 +1339,13 @@ namespace dlib { namespace tt const tensor& gradient_input ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::gelu_gradient(grad, src, gradient_input); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::gelu_gradient(grad, src, gradient_input); + ) } // ---------------------------------------------------------------------------------------- @@ -1308,12 +1357,14 @@ namespace dlib { namespace tt ) { DLIB_CASSERT(beta > 0); -#ifdef DLIB_USE_CUDA - if (use_cuda()) + + IF_DLIB_USE_CUDA( cuda::smelu(dest, src, beta); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::smelu(dest, src, beta); + ) } void smelu_gradient ( @@ -1324,12 +1375,14 @@ namespace dlib { namespace tt ) { DLIB_CASSERT(beta > 0); -#ifdef DLIB_USE_CUDA - if (use_cuda()) + + IF_DLIB_USE_CUDA( cuda::smelu_gradient(grad, dest, gradient_input, beta); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::smelu_gradient(grad, dest, gradient_input, beta); + ) } // ---------------------------------------------------------------------------------------- @@ -1339,12 +1392,13 @@ namespace dlib { namespace tt const tensor& src ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::silu(dest,src); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::silu(dest,src); + ) } void silu_gradient ( @@ -1353,12 +1407,13 @@ namespace dlib { namespace tt const tensor& gradient_input ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::silu_gradient(grad, src, gradient_input); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::silu_gradient(grad, src, gradient_input); + ) } // ---------------------------------------------------------------------------------------- @@ -1372,12 +1427,13 @@ namespace dlib { namespace tt long src_channel_stride ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::resize_bilinear(dest,dest_row_stride,dest_channel_stride, src,src_row_stride,src_channel_stride); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::resize_bilinear(dest,dest_row_stride,dest_channel_stride, src,src_row_stride,src_channel_stride); + ) } void resize_bilinear_gradient ( @@ -1389,12 +1445,13 @@ namespace dlib { namespace tt long gradient_input_channel_stride ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::resize_bilinear_gradient(grad,grad_row_stride,grad_channel_stride, gradient_input,gradient_input_row_stride,gradient_input_channel_stride); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::resize_bilinear_gradient(grad,grad_row_stride,grad_channel_stride, gradient_input,gradient_input_row_stride,gradient_input_channel_stride); + ) } // ------------------------------------------------------------------------------------ @@ -1407,12 +1464,13 @@ namespace dlib { namespace tt const tensor& src ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::reorg(add_to, dest, row_stride, col_stride, src); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::reorg(add_to, dest, row_stride, col_stride, src); + ) } void reorg_gradient ( @@ -1423,12 +1481,13 @@ namespace dlib { namespace tt const tensor& gradient_input ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::reorg_gradient(add_to, grad, row_stride, col_stride, gradient_input); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::reorg_gradient(add_to, grad, row_stride, col_stride, gradient_input); + ) } // ------------------------------------------------------------------------------------ @@ -1442,12 +1501,13 @@ namespace dlib { namespace tt size_t count_k ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::copy_tensor(add_to, dest, dest_k_offset, src, src_k_offset, count_k); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::copy_tensor(add_to, dest, dest_k_offset, src, src_k_offset, count_k); + ) } // ---------------------------------------------------------------------------------------- @@ -1461,12 +1521,13 @@ namespace dlib { namespace tt size_t k, size_t nr, size_t nc ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::copy_tensor(add_to, dest, dk, dnr, dnc , src, sk, snr, snc, k, nr, nc); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::copy_tensor(add_to, dest, dk, dnr, dnc, src, sk, snr, snc, k, nr, nc); + ) } // ---------------------------------------------------------------------------------------- @@ -1477,12 +1538,13 @@ namespace dlib { namespace tt resizable_tensor& out ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( finv(m,out); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( out = dlib::inv(mat(m)); + ) } // ---------------------------------------------------------------------------------------- @@ -1493,12 +1555,13 @@ namespace dlib { namespace tt const tensor& src ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::transpose(add_to, dest, src); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::transpose(add_to, dest, src); + ) } // ---------------------------------------------------------------------------------------- @@ -1509,12 +1572,13 @@ namespace dlib { namespace tt const tensor& embs ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::embeddings(dest, src, embs); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::embeddings(dest, src, embs); + ) } void embeddings_gradient( @@ -1526,12 +1590,13 @@ namespace dlib { namespace tt bool scale ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda::embeddings_gradient(prev, gradient_input, grads, freqs, learning_rate, scale); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu::embeddings_gradient(prev, gradient_input, grads, freqs, learning_rate, scale); + ) } // ---------------------------------------------------------------------------------------- diff --git a/dlib/cuda/tensor_tools.h b/dlib/cuda/tensor_tools.h index 06a7fc01da..f9ec1d11fe 100644 --- a/dlib/cuda/tensor_tools.h +++ b/dlib/cuda/tensor_tools.h @@ -15,6 +15,18 @@ #include "../geometry/rectangle.h" #include "../test_for_odr_violations.h" +#ifdef DLIB_USE_CUDA +#define IF_DLIB_USE_CUDA(...) if (use_cuda()) { __VA_ARGS__ } +#else +#define IF_DLIB_USE_CUDA(...) +#endif + +#ifdef DLIB_USE_CUDA +#define IF_DLIB_NOT_USE_CUDA(...) if (!use_cuda()) { __VA_ARGS__ } +#else +#define IF_DLIB_NOT_USE_CUDA(...) __VA_ARGS__ +#endif + namespace dlib { bool dnn_prefer_fastest_algorithms(); @@ -1078,12 +1090,13 @@ namespace dlib { namespace tt void clear( ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda_impl.clear(); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu_impl.clear(); + ) } void operator() ( @@ -1093,12 +1106,13 @@ namespace dlib { namespace tt const tensor& filters ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda_impl(add_to_output,output,data,filters); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu_impl(add_to_output,output,data,filters); + ) } /*! requires @@ -1127,12 +1141,13 @@ namespace dlib { namespace tt const tensor& filters ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda_impl(add_to_output,output,data,filters); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu_impl(add_to_output,output,data,filters); + ) } /*! requires @@ -1163,12 +1178,13 @@ namespace dlib { namespace tt bool use_relu ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda_impl(add_to_output,output,data,filters,biases,use_relu); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu_impl(add_to_output,output,data,filters,biases,use_relu); + ) } /*! requires @@ -1203,12 +1219,13 @@ namespace dlib { namespace tt bool use_relu ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda_impl(add_to_output,output,data,filters,biases,use_relu); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu_impl(add_to_output,output,data,filters,biases,use_relu); + ) } /*! requires @@ -1239,12 +1256,13 @@ namespace dlib { namespace tt tensor& data_gradient ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda_impl.get_gradient_for_data(add_to_output,gradient_input,filters,data_gradient); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu_impl.get_gradient_for_data(add_to_output,gradient_input,filters,data_gradient); + ) } /*! requires @@ -1282,12 +1300,13 @@ namespace dlib { namespace tt tensor& filters_gradient ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda_impl.get_gradient_for_filters(add_to_output,gradient_input,data,filters_gradient); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu_impl.get_gradient_for_filters(add_to_output,gradient_input,data,filters_gradient); + ) } /*! requires @@ -1328,12 +1347,13 @@ namespace dlib { namespace tt int padding_x ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda_impl.setup(data,filters,stride_y,stride_x,padding_y,padding_x); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu_impl.setup(data,filters,stride_y,stride_x,padding_y,padding_x); + ) } /*! requires @@ -1383,12 +1403,13 @@ namespace dlib { namespace tt void clear( ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda_impl.clear(); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu_impl.clear(); + ) } void setup_max_pooling( @@ -1400,12 +1421,13 @@ namespace dlib { namespace tt int padding_x ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda_impl.setup_max_pooling(window_height, window_width, stride_y, stride_x, padding_y, padding_x); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu_impl.setup_max_pooling(window_height, window_width, stride_y, stride_x, padding_y, padding_x); + ) } /*! requires @@ -1429,12 +1451,13 @@ namespace dlib { namespace tt int padding_x ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda_impl.setup_avg_pooling(window_height, window_width, stride_y, stride_x, padding_y, padding_x); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu_impl.setup_avg_pooling(window_height, window_width, stride_y, stride_x, padding_y, padding_x); + ) } /*! requires @@ -1452,12 +1475,13 @@ namespace dlib { namespace tt bool does_max_pooling( ) const { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( return cuda_impl.does_max_pooling(); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( return cpu_impl.does_max_pooling(); + ) } void operator() ( @@ -1465,12 +1489,13 @@ namespace dlib { namespace tt const tensor& src ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda_impl(dest, src); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu_impl(dest, src); + ) } /*! requires @@ -1501,12 +1526,13 @@ namespace dlib { namespace tt tensor& grad ) { -#ifdef DLIB_USE_CUDA - if (use_cuda()) + IF_DLIB_USE_CUDA( cuda_impl.get_gradient(gradient_input, dest, src, grad); - else -#endif + ) + + IF_DLIB_NOT_USE_CUDA( cpu_impl.get_gradient(gradient_input, dest, src, grad); + ) } /*! requires From dc33fa6ab3e9e8cda1732f3aed49230ec828d19b Mon Sep 17 00:00:00 2001 From: Kane Scipioni Date: Tue, 6 May 2025 08:17:49 -0500 Subject: [PATCH 17/39] Added docs for new public functions --- dlib/dnn/core_abstract.h | 18 ++++++++++++++++++ 1 file changed, 18 insertions(+) diff --git a/dlib/dnn/core_abstract.h b/dlib/dnn/core_abstract.h index 491183a685..70f1417e11 100644 --- a/dlib/dnn/core_abstract.h +++ b/dlib/dnn/core_abstract.h @@ -196,6 +196,24 @@ namespace dlib - #dnn_prefer_fastest_algorithms() == false !*/ + bool use_cuda( + ); + /*! + ensures + - If dlib should use the CUDA implementation of a deep neural network + then this function returns true and false otherwise. + - On program startup this function will return true if DLIB_USE_CUDA is defined. + - This function always returns false if DLIB_USE_CUDA is not defined. + !*/ + + void set_use_cuda( + bool flag + ); + /*! + ensures + - #use_cuda() == flag + !*/ + // ---------------------------------------------------------------------------------------- template < From 652af017821152c418d6c1dee56f93127583e028 Mon Sep 17 00:00:00 2001 From: Cydral <53169060+Cydral@users.noreply.github.com> Date: Sat, 3 May 2025 17:17:38 +0200 Subject: [PATCH 18/39] Add linear_ layer for neural networks (#3074) --- dlib/dnn/layers.h | 200 ++++++++++++++++++++++++++++ dlib/dnn/layers_abstract.h | 260 +++++++++++++++++++++++++++++++++++++ dlib/test/dnn.cpp | 75 +++++++++++ 3 files changed, 535 insertions(+) diff --git a/dlib/dnn/layers.h b/dlib/dnn/layers.h index 0a0c547f33..123d49f523 100644 --- a/dlib/dnn/layers.h +++ b/dlib/dnn/layers.h @@ -2143,6 +2143,206 @@ namespace dlib > using fc_no_bias = add_layer, SUBNET>; +// ---------------------------------------------------------------------------------------- + + enum linear_bias_mode { LINEAR_HAS_BIAS = 0, LINEAR_NO_BIAS = 1 }; + + template < + unsigned long num_outputs_, + linear_bias_mode bias_mode_ + > + class linear_ + { + static_assert(num_outputs_ > 0, "The number of outputs from a linear_ layer must be > 0"); + + public: + linear_() : + num_outputs(num_outputs_), + num_inputs(0), + learning_rate_multiplier(1), + bias_mode(bias_mode_) { + } + + double get_learning_rate_multiplier() const { return learning_rate_multiplier; } + void set_learning_rate_multiplier(double val) { learning_rate_multiplier = val; } + + unsigned long get_num_inputs() const { return num_inputs; } + unsigned long get_num_outputs() const { return num_outputs; } + void set_num_outputs(long num) + { + DLIB_CASSERT(num > 0, "The number of outputs must be > 0, but num == " << num); + if (num != (long)num_outputs) + { + DLIB_CASSERT(get_layer_params().size() == 0, + "You can't change the number of filters in linear_ if the parameter tensor has already been allocated."); + num_outputs = num; + } + } + linear_bias_mode get_bias_mode() const { return bias_mode; } + + template + void setup(const SUBNET& sub) + { + num_inputs = sub.get_output().nc(); + if (bias_mode == LINEAR_HAS_BIAS) + params.set_size(num_inputs + 1, num_outputs); + else + params.set_size(num_inputs, num_outputs); + + dlib::rand rnd(std::rand()); + randomize_parameters(params, num_inputs + num_outputs, rnd); + weights = alias_tensor(num_inputs, num_outputs); + + if (bias_mode == LINEAR_HAS_BIAS) { + biases = alias_tensor(1, num_outputs); + biases(params, weights.size()) = 0; + } + } + + template + void forward(const SUBNET& sub, resizable_tensor& output) + { + const auto& prev_output = sub.get_output(); + DLIB_CASSERT((long)num_inputs == prev_output.nc(), + "The size of the input tensor to this linear layer doesn't match the size the linear layer was trained with."); + output.set_size(prev_output.num_samples(), prev_output.k(), prev_output.nr(), num_outputs); + + auto o = alias_tensor(output.num_samples() * output.k() * output.nr(), num_outputs)(output, 0); + auto so = alias_tensor(prev_output.num_samples() * prev_output.k() * prev_output.nr(), num_inputs)(prev_output, 0); + + auto w = weights(params, 0); + tt::gemm(0, (tensor&)o, 1, so, false, w, false); + + if (bias_mode == LINEAR_HAS_BIAS) + { + auto b = biases(params, weights.size()); + tt::add(1, (tensor&)o, 1, b); + } + } + + template + void backward(const tensor& gradient_input, SUBNET& sub, tensor& params_grad) + { + auto gi = alias_tensor(gradient_input.num_samples() * gradient_input.k() * gradient_input.nr(), num_outputs)(gradient_input, 0); + if (learning_rate_multiplier != 0) + { + const auto& prev_output = sub.get_output(); + auto pw = weights(params_grad, 0); + auto so = alias_tensor(prev_output.num_samples() * prev_output.k() * prev_output.nr(), num_inputs)(prev_output, 0); + tt::gemm(0, pw, learning_rate_multiplier, so, true, gi, false); + + if (bias_mode == LINEAR_HAS_BIAS) + { + auto pb = biases(params_grad, weights.size()); + tt::assign_bias_gradient(pb, gi); + } + } + + const auto& prev_gradient = sub.get_gradient_input(); + auto sgi = alias_tensor(prev_gradient.num_samples() * prev_gradient.k() * prev_gradient.nr(), num_inputs)(prev_gradient, 0); + auto w = weights(params, 0); + tt::gemm(1, (tensor&)sgi, 1, gi, false, w, true); + } + + alias_tensor_instance get_weights() { return weights(params, 0); } + alias_tensor_const_instance get_weights() const { return weights(params, 0); } + alias_tensor_instance get_biases() + { + static_assert(bias_mode == LINEAR_HAS_BIAS, "This linear_ layer doesn't have a bias vector " + "to be retrieved, as per template parameter 'bias_mode'."); + return biases(params, weights.size()); + } + alias_tensor_const_instance get_biases() const + { + static_assert(bias_mode == LINEAR_HAS_BIAS, "This linear_ layer doesn't have a bias vector " + "to be retrieved, as per template parameter 'bias_mode'."); + return biases(params, weights.size()); + } + + inline dpoint map_input_to_output(const dpoint& p) const { return p; } + inline dpoint map_output_to_input(const dpoint& p) const { return p; } + + const tensor& get_layer_params() const { return params; } + tensor& get_layer_params() { return params; } + + friend void serialize(const linear_& item, std::ostream& out) + { + serialize("linear_", out); + serialize(item.num_outputs, out); + serialize(item.num_inputs, out); + serialize(item.params, out); + serialize(item.weights, out); + serialize(item.biases, out); + serialize((int)item.bias_mode, out); + serialize(item.learning_rate_multiplier, out); + } + + friend void deserialize(linear_& item, std::istream& in) + { + std::string version; + deserialize(version, in); + if (version == "linear_") + { + deserialize(item.num_outputs, in); + deserialize(item.num_inputs, in); + deserialize(item.params, in); + deserialize(item.weights, in); + deserialize(item.biases, in); + int bmode; + deserialize(bmode, in); + item.bias_mode = static_cast(bmode); + if (bias_mode_ != item.bias_mode) throw serialization_error("Wrong bias_mode found while deserializing dlib::linear_"); + deserialize(item.learning_rate_multiplier, in); + } + else + { + throw serialization_error("Unexpected version '" + version + "' found while deserializing dlib::linear_."); + } + } + + friend std::ostream& operator<<(std::ostream& out, const linear_& item) + { + out << "linear\t (num_outputs=" << item.num_outputs; + if (item.bias_mode == LINEAR_HAS_BIAS) + out << ", bias=true"; + else + out << ", bias=false"; + out << ")"; + out << " learning_rate_mult=" << item.learning_rate_multiplier; + return out; + } + + friend void to_xml(const linear_& item, std::ostream& out) + { + out << "\n"; + out << mat(item.params); + out << "\n"; + } + + private: + unsigned long num_inputs; + unsigned long num_outputs; + double learning_rate_multiplier; + linear_bias_mode bias_mode; + resizable_tensor params; + alias_tensor weights, biases; + }; + + template < + unsigned long num_outputs, + typename SUBNET + > + using linear = add_layer, SUBNET>; + + template < + unsigned long num_outputs, + typename SUBNET + > + using linear_no_bias = add_layer, SUBNET>; + // ---------------------------------------------------------------------------------------- class dropout_ diff --git a/dlib/dnn/layers_abstract.h b/dlib/dnn/layers_abstract.h index ef2de8e6fe..e5f2d340e0 100644 --- a/dlib/dnn/layers_abstract.h +++ b/dlib/dnn/layers_abstract.h @@ -689,6 +689,266 @@ namespace dlib > using fc_no_bias = add_layer, SUBNET>; + // ---------------------------------------------------------------------------------------- + +// ---------------------------------------------------------------------------------------- + + enum linear_bias_mode + { + LINEAR_HAS_BIAS, + LINEAR_NO_BIAS + }; + + template < + unsigned long num_outputs, + linear_bias_mode bias_mode + > + class linear_ + { + /*! + REQUIREMENTS ON num_outputs + num_outputs > 0 + + WHAT THIS OBJECT REPRESENTS + This is an implementation of a linear layer, which applies a linear + transformation to the input data. For a layer with bias, the transformation + is: + output = input * weights + bias + For a layer without bias, it's simply: + output = input * weights + + The input tensor can have any number of sample, k (channel), and nr (row) + dimensions, but the nc (column) dimension must match the number of input features. + The output tensor will have the same dimensions as the input tensor, except for + the nc dimension which will be equal to num_outputs. + + This layer is similar to the fc_ layer, but optimized for the case where the + input and output tensors maintain the same dimensions, excluding the feature + dimension (nc). This makes it useful for working with multi-dimensional data. + !*/ + + public: + linear_( + ); + /*! + ensures + - #get_num_outputs() == num_outputs + - #get_bias_mode() == bias_mode + - #get_learning_rate_multiplier() == 1 + !*/ + + double get_learning_rate_multiplier( + ) const; + /*! + ensures + - returns a multiplier that will be applied to the gradient of this layer during + training. This value appears as a multiplicative factor in the update rule. So + if get_learning_rate_multiplier() == 1 then the learning rate will be multiplied + by 1 and thus not modified. However, if get_learning_rate_multiplier() == 0.1 then + the learning rate will be multiplied by 0.1, making the layer update 10 times + slower than it would otherwise be. + !*/ + + void set_learning_rate_multiplier( + double val + ); + /*! + ensures + - #get_learning_rate_multiplier() == val + !*/ + + unsigned long get_num_inputs( + ) const; + /*! + ensures + - Returns the number of input features this layer expects. + - For an uninitialized layer (i.e., one that has not seen any data during setup + or forward pass), this will be zero. + !*/ + + unsigned long get_num_outputs( + ) const; + /*! + ensures + - Returns the number of output features this layer produces. + I.e., this value is num_outputs. + !*/ + + void set_num_outputs( + long num + ); + /*! + requires + - num > 0 + ensures + - #get_num_outputs() == num + throws + - std::runtime_error if this function is called after the layer parameters + have been allocated and the new number of outputs doesn't match the + previously set number of outputs. + !*/ + + linear_bias_mode get_bias_mode( + ) const; + /*! + ensures + - Returns a value indicating whether this layer has a bias term. + I.e. returns bias_mode. + !*/ + + template + void setup( + const SUBNET& sub + ); + /*! + ensures + - Performs the necessary setup work to process data through this layer. + - Sets the input size based on the dimensions of the input tensor from sub. + - Allocates the parameter tensor and initializes its values. + - #get_num_inputs() == the number of columns in sub.get_output() (i.e., nc). + !*/ + + template + void forward( + const SUBNET& sub, + resizable_tensor& output + ); + /*! + requires + - setup() has been called + - sub.get_output().nc() == get_num_inputs() + ensures + - Applies the linear transformation to the input tensor from sub and stores + the results in output. + - #output.num_samples() == sub.get_output().num_samples() + - #output.k() == sub.get_output().k() + - #output.nr() == sub.get_output().nr() + - #output.nc() == get_num_outputs() + !*/ + + template + void backward( + const tensor& gradient_input, + SUBNET& sub, + tensor& params_grad + ); + /*! + requires + - setup() has been called + - sub.get_output().nc() == get_num_inputs() + - gradient_input has the same dimensions as the output of forward() + ensures + - Computes the gradients of this layer with respect to the parameters + and the input tensor, and updates the corresponding gradient tensors. + - Updates params_grad based on the gradients of the weights + and biases (if present). + - Updates sub's gradient_input based on the gradients of the + inputs to this layer. + !*/ + + alias_tensor_instance get_weights( + ); + /*! + requires + - setup() has been called + ensures + - Returns a reference to the weights matrix of this layer. + !*/ + + alias_tensor_const_instance get_weights( + ) const; + /*! + requires + - setup() has been called + ensures + - Returns a const reference to the weights matrix of this layer. + !*/ + + alias_tensor_instance get_biases( + ); + /*! + requires + - bias_mode == LINEAR_HAS_BIAS + - setup() has been called + ensures + - Returns a reference to the bias vector of this layer. + throws + - static_assert failure if bias_mode != LINEAR_HAS_BIAS + !*/ + + alias_tensor_const_instance get_biases( + ) const; + /*! + requires + - bias_mode == LINEAR_HAS_BIAS + - setup() has been called + ensures + - Returns a const reference to the bias vector of this layer. + throws + - static_assert failure if bias_mode != LINEAR_HAS_BIAS + !*/ + + dpoint map_input_to_output( + const dpoint& p + ) const; + /*! + ensures + - Returns p, since the linear layer maintains the same spatial dimensions. + !*/ + + dpoint map_output_to_input( + const dpoint& p + ) const; + /*! + ensures + - Returns p, since the linear layer maintains the same spatial dimensions. + !*/ + + const tensor& get_layer_params( + ) const; + /*! + ensures + - Returns the parameters that define this layer, i.e., the weights and biases + (if present) that are updated during training. + !*/ + + tensor& get_layer_params( + ); + /*! + ensures + - Returns the parameters that define this layer, i.e., the weights and biases + (if present) that are updated during training. + !*/ + + friend void serialize(const linear_& item, std::ostream& out); + friend void deserialize(linear_& item, std::istream& in); + /*! + provides serialization support + !*/ + }; + + template < + unsigned long num_outputs, + typename SUBNET + > + using linear = add_layer, SUBNET>; + /*! + This is a layer that applies a linear transformation with bias to the input: + output = input * weights + bias + !*/ + + template < + unsigned long num_outputs, + typename SUBNET + > + using linear_no_bias = add_layer, SUBNET>; + /*! + This is a layer that applies a linear transformation without bias to the input: + output = input * weights + !*/ + + // ---------------------------------------------------------------------------------------- + // ---------------------------------------------------------------------------------------- struct num_con_outputs diff --git a/dlib/test/dnn.cpp b/dlib/test/dnn.cpp index 1da3bac777..de26a1e1d3 100644 --- a/dlib/test/dnn.cpp +++ b/dlib/test/dnn.cpp @@ -2419,6 +2419,24 @@ void test_embeddings() auto res = test_layer(l); DLIB_TEST_MSG(res, res); } + { + print_spinner(); + linear_<1, LINEAR_NO_BIAS> l; + auto res = test_layer(l); + DLIB_TEST_MSG(res, res); + } + { + print_spinner(); + linear_<5, LINEAR_NO_BIAS> l; + auto res = test_layer(l); + DLIB_TEST_MSG(res, res); + } + { + print_spinner(); + linear_<4, LINEAR_NO_BIAS> l; + auto res = test_layer(l); + DLIB_TEST_MSG(res, res); + } { print_spinner(); relu_ l; @@ -3527,6 +3545,62 @@ void test_multm_prev() DLIB_TEST_MSG(error_after < 1e-6, "Autoencoder error after training = " << error_after); } +// ---------------------------------------------------------------------------------------- + void test_linear() + { + print_spinner(); + + // Define the network + cout << "ICI !!!" << endl; + using net_type = tag2>>>>; + net_type net; + + // Input tensor + const int n_samples = 3, k = 1; + std::vector> x(n_samples); + matrix xtmp(2, 4); + xtmp = 1.0f, 2.0f, 3.0f, 4.0f, + 5.0f, 6.0f, 7.0f, 8.0f; + x[0] = xtmp; + xtmp = 9.0f, 10.0f, 11.0f, 12.0f, + 13.0f, 14.0f, 15.0f, 16.0f; + x[1] = xtmp; + xtmp = 17.0f, 18.0f, 19.0f, 20.0f, + 21.0f, 22.0f, 23.0f, 24.0f; + x[2] = xtmp; + + // Convert input matrix to tensor + resizable_tensor input_tensor; + net.to_tensor(&x[0], &x[0] + n_samples, input_tensor); + net.forward(input_tensor); + + // Get the internal linear weights + matrix w = mat(layer(net).subnet().layer_details().get_weights()); + + // Theoretical calculation of the output + std::vector> expected_outputs(n_samples); + for (int i = 0; i < n_samples; ++i) { + matrix input_matrix = x[i]; + expected_outputs[i] = input_matrix * w; + } + + // Compare output tensor with expected output + auto& net_output = layer(net).get_output(); + + // Display results + for (int i = 0; i < n_samples; ++i) { + matrix output_sample; + output_sample.set_size(2, 6); + for (long r = 0; r < output_sample.nr(); ++r) { + for (long c = 0; c < output_sample.nc(); ++c) { + output_sample(r, c) = net_output.host()[tensor_index(net_output, i, 0, r, c)]; + } + } + DLIB_TEST_MSG(max(abs(output_sample - expected_outputs[i])) < 1e-5, + "linear layer - sample " + std::to_string(i)); + } + } + // ---------------------------------------------------------------------------------------- void test_loss_mean_squared_per_channel_and_pixel() @@ -5109,6 +5183,7 @@ void test_multm_prev() test_simple_linear_regression_with_mult_prev(); test_multioutput_linear_regression(); test_simple_autoencoder(); + test_linear(); test_loss_mean_squared_per_channel_and_pixel(); test_loss_binary_log_per_pixel_learned_params_on_trivial_two_pixel_task(); test_loss_binary_log_per_pixel_outputs_on_trivial_task(); From f97db8fc3af0a54fc9906fcffca49bbdf37ab6b5 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Adri=C3=A0?= <1671644+arrufat@users.noreply.github.com> Date: Sat, 10 May 2025 08:30:13 +0900 Subject: [PATCH 19/39] ci: remove unsupported ubuntu 20.04 (#3075) --- .github/workflows/build_cpp.yml | 54 --------------------------------- 1 file changed, 54 deletions(-) diff --git a/.github/workflows/build_cpp.yml b/.github/workflows/build_cpp.yml index 720eb4defd..bbbd24449b 100644 --- a/.github/workflows/build_cpp.yml +++ b/.github/workflows/build_cpp.yml @@ -121,60 +121,6 @@ jobs: - name: Test BLAS bindings run: build_blas_bindings/dtest --runall -q - ubuntu-20_04-gcc-7: - runs-on: 'ubuntu-20.04' - steps: - - uses: actions/checkout@v2 - - - name: Install dependencies - run: | - sudo apt update - sudo apt install libwebp-dev make yasm - - - name: Install gcc 7 - run: | - sudo apt install gcc-7 g++-7 - - - name: Cache FFmpeg 3.2.18 - uses: actions/cache@v3 - id: cache-ffmpeg3 - with: - path: /home/runner/ffmpeg-n3.2.18_installation - key: ffmpeg-n3.2.18_try1 - - - name: Build FFmpeg 3.2.18 - if: steps.cache-ffmpeg3.outputs.cache-hit != 'true' - run: | - wget https://github.com/FFmpeg/FFmpeg/archive/refs/tags/n3.2.18.tar.gz - tar -xf n3.2.18.tar.gz - cd FFmpeg-n3.2.18 - ./configure --prefix=/home/runner/ffmpeg-n3.2.18_installation --disable-doc --disable-programs - make -j4 - make install - cd .. - - - name: Configure - run: | - export CC=/usr/bin/gcc-7 - export CXX=/usr/bin/g++-7 - cmake ${{ github.workspace }}/dlib/test -B build -DCMAKE_PREFIX_PATH=/home/runner/ffmpeg-n3.2.18_installation - - - name: Build just tests - run: cmake --build build --config Release --target dtest --parallel 4 - - - name: Test - run: build/dtest --runall -q - - - name: Build examples, etc - run: cmake --build build --config Release --parallel 2 - - # Test cmake scrips can build standalone dlib as a shared library - - name: Configure dlib as shared library - run: cmake ${{ github.workspace }}/dlib -B build_dlib_shared -DBUILD_SHARED_LIBS=1 -DCMAKE_PREFIX_PATH=/home/runner/ffmpeg-n3.2.18_installation - - - name: Build dlib as shared library - run: cmake --build build_dlib_shared --parallel 4 - ubuntu-latest-clang-default-avx: runs-on: 'ubuntu-latest' steps: From ec1881bdd2fae57fa5cb29212bf17ab2dcbc2148 Mon Sep 17 00:00:00 2001 From: Cydral <53169060+Cydral@users.noreply.github.com> Date: Fri, 23 May 2025 13:27:34 +0200 Subject: [PATCH 20/39] Add `reshape_to` layer for flexible tensor reshaping/rescaling (#3076) * Implementation of linear_ layer for neural networks. This layer provides an optimized linear transformation for multi-dimensional inputs. * Minor change * Update dlib/dnn/layers.h Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * Add reshape_to and flatten layers to Dlib's DNN module * Missing update to "visitors.h" * format fixing for reshape_to * Update dlib/test/dnn.cpp --------- Co-authored-by: Davis E. King Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> --- dlib/dnn/layers.h | 180 +++++++++++++++++++++++++++++++++++++ dlib/dnn/layers_abstract.h | 170 +++++++++++++++++++++++++++++++++++ dlib/dnn/visitors.h | 23 +++++ dlib/test/dnn.cpp | 48 +++++++++- 4 files changed, 420 insertions(+), 1 deletion(-) diff --git a/dlib/dnn/layers.h b/dlib/dnn/layers.h index 123d49f523..7ec8b1a956 100644 --- a/dlib/dnn/layers.h +++ b/dlib/dnn/layers.h @@ -975,6 +975,186 @@ namespace dlib > using resize_to = add_layer, SUBNET>; +// ---------------------------------------------------------------------------------------- + + template + class reshape_to_ + { + public: + explicit reshape_to_() : + output_k(k_), + output_nr(nr_), + output_nc(nc_) + { + static_assert(k_ == -1 || k_ > 0, "Output k must be positive or -1"); + static_assert(nr_ == -1 || nr_ > 0, "Output nr must be positive or -1"); + static_assert(nc_ == -1 || nc_ > 0, "Output nc must be positive or -1"); + + input_k = input_nr = input_nc = 0; + needs_rescale = false; + } + + // Getters for dimensions + long get_output_k() const { return output_k; } + long get_output_nr() const { return output_nr; } + long get_output_nc() const { return output_nc; } + + // Setters for dimensions + void set_output_k(long k) { + DLIB_CASSERT(k == -1 || k > 0, "Output k must be positive or -1 to keep original dimension"); + output_k = k; + } + void set_output_nr(long nr) { + DLIB_CASSERT(nr == -1 || nr > 0, "output nr must be positive or -1 to keep original dimension"); + output_nr = nr; + } + void set_output_nc(long nc) { + DLIB_CASSERT(nc == -1 || nc > 0, "output nc must be positive or -1 to keep original dimension"); + output_nc = nc; + } + + template + void setup(const SUBNET& sub) + { + const auto& input = sub.get_output(); + input_k = input.k(); + input_nr = input.nr(); + input_nc = input.nc(); + + // Calculate output dimensions using input dims where target is -1 + if (k_ == -1) output_k = input_k; + if (nr_ == -1) output_nr = input_nr; + if (nc_ == -1) output_nc = input_nc; + + // Check if this is well a pure reshape + long input_elements = input_k * input_nr * input_nc; + long output_elements = output_k * output_nr * output_nc; + if (input_elements != output_elements && input_k == output_k) needs_rescale = true; + DLIB_CASSERT(input_elements == output_elements || needs_rescale, + "Cannot reshape tensor of " << input_elements << + " elements into shape with " << output_elements << " elements. " << + "For spatial rescaling, the channel dimension (k) must remain constant."); + } + + template + void forward(const SUBNET& sub, resizable_tensor& output) + { + // Set the output size (always preserving batch dimension) + const tensor& input = sub.get_output(); + output.set_size(input.num_samples(), output_k, output_nr, output_nc); + + if (!needs_rescale) + { + // Create an alias of the input tensor with the output shape + alias_tensor input_alias(output.num_samples(), output_k, output_nr, output_nc); + // Get a view of the input tensor with the new shape + auto input_reshaped = input_alias(const_cast(input), 0); + // Copy the view to the output tensor + tt::copy_tensor(false, output, 0, input_reshaped, 0, input_reshaped.k()); + } + else + { + // Only spatial dimensions need to be resized + tt::resize_bilinear(output, input); + } + } + + template + void backward(const tensor& gradient_input, SUBNET& sub, tensor& /*params_grad*/) + { + auto& grad = sub.get_gradient_input(); + + if (!needs_rescale) { + // Create an alias of the gradient tensor with the original input shape + alias_tensor grad_alias(grad.num_samples(), grad.k(), grad.nr(), grad.nc()); + // Get a view of the input gradient with the required shape + auto grad_reshaped = grad_alias(const_cast(gradient_input), 0); + // Copy the view to the output gradient + tt::copy_tensor(true, grad, 0, grad_reshaped, 0, grad_reshaped.k()); + } + else + { + // Only spatial dimensions were resized + tt::resize_bilinear_gradient(grad, gradient_input); + } + } + + // Mapping functions for coordinate transformations + inline dpoint map_input_to_output(const dpoint& p) const { + double scale_x = output_nc / static_cast(input_nc); + double scale_y = output_nr / static_cast(input_nr); + return dpoint(p.x() * scale_x, p.y() * scale_y); + } + inline dpoint map_output_to_input(const dpoint& p) const { + double scale_x = input_nc / static_cast(output_nc); + double scale_y = input_nr / static_cast(output_nr); + return dpoint(p.x() * scale_x, p.y() * scale_y); + } + + const tensor& get_layer_params() const { return params; } + tensor& get_layer_params() { return params; } + + friend void serialize(const reshape_to_& item, std::ostream& out) + { + serialize("reshape_to_", out); + serialize(item.input_k, out); + serialize(item.input_nr, out); + serialize(item.input_nc, out); + serialize(item.output_k, out); + serialize(item.output_nr, out); + serialize(item.output_nc, out); + serialize(item.needs_rescale, out); + } + + friend void deserialize(reshape_to_& item, std::istream& in) + { + std::string version; + deserialize(version, in); + if (version != "reshape_to_") + throw serialization_error("Unexpected version '" + version + "' found while deserializing dlib::reshape_to_."); + deserialize(item.input_k, in); + deserialize(item.input_nr, in); + deserialize(item.input_nc, in); + deserialize(item.output_k, in); + deserialize(item.output_nr, in); + deserialize(item.output_nc, in); + deserialize(item.needs_rescale, in); + } + + friend std::ostream& operator<<(std::ostream& out, const reshape_to_& item) + { + out << "reshape_to ("; + out << "k=" << std::to_string(item.output_k); + out << ", nr=" << std::to_string(item.output_nr); + out << ", nc=" << std::to_string(item.output_nc); + out << ", mode=" << (item.needs_rescale ? "spatial_rescale" : "pure_reshape"); + out << ")"; + return out; + } + + friend void to_xml(const reshape_to_& item, std::ostream& out) + { + out << "\n"; + } + + private: + long input_k, input_nr, input_nc; // Input dimensions + long output_k, output_nr, output_nc; // Output dimensions + bool needs_rescale; + resizable_tensor params; // No trainable parameters + }; + + template + using reshape_to = add_layer, SUBNET>; + + template + using flatten = add_layer, SUBNET>; + // ---------------------------------------------------------------------------------------- template < diff --git a/dlib/dnn/layers_abstract.h b/dlib/dnn/layers_abstract.h index e5f2d340e0..f0512d7d4c 100644 --- a/dlib/dnn/layers_abstract.h +++ b/dlib/dnn/layers_abstract.h @@ -1642,6 +1642,176 @@ namespace dlib > using resize_to = add_layer, SUBNET>; +// ---------------------------------------------------------------------------------------- + + template + class reshape_to_ + { + /*! + REQUIREMENTS ON TEMPLATE ARGUMENTS + - k_, nr_, and nc_ must be either -1 or greater than 0. + + WHAT THIS OBJECT REPRESENTS + This is an implementation of the EXAMPLE_COMPUTATIONAL_LAYER_ interface + defined above. It defines a layer that reshapes or resizes an input tensor + into a different shape. The layer operates in two modes: + + 1. Pure Reshape Mode: When the total number of elements in the input tensor + equals the total number of elements in the output tensor, this layer + performs a simple reshaping operation without changing the values. + + 2. Spatial Rescaling Mode: When the channel dimension (k) remains constant + but the total number of elements changes, this layer performs bilinear + interpolation to resize the spatial dimensions while preserving the + channel information. + + The dimensions of the output tensor are determined by the template parameters: + - If k_ is -1, the output tensor will have the same number of channels as the input. + - If nr_ is -1, the output tensor will have the same number of rows as the input. + - If nc_ is -1, the output tensor will have the same number of columns as the input. + + Setting a value of -1 for any dimension means "keep the original dimension from the input." + + Note that this layer will throw an exception if you attempt to change both the + channel count (k) and the total number of elements. Either: + - Keep the total number of elements the same (Pure Reshape Mode), or + - Keep the channel count the same and only change spatial dimensions (Spatial Rescaling Mode) + !*/ + + public: + explicit reshape_to_(); + /*! + ensures + - #get_output_k() == k_ + - #get_output_nr() == nr_ + - #get_output_nc() == nc_ + !*/ + + long get_output_k() const; + /*! + ensures + - Returns the number of channels in the output tensor. If this value is -1, + then the output will have the same number of channels as the input. + !*/ + + long get_output_nr() const; + /*! + ensures + - Returns the number of rows in the output tensor. If this value is -1, + then the output will have the same number of rows as the input. + !*/ + + long get_output_nc() const; + /*! + ensures + - Returns the number of columns in the output tensor. If this value is -1, + then the output will have the same number of columns as the input. + !*/ + + void set_output_k(long k); + /*! + requires + - k == -1 || k > 0 + ensures + - #get_output_k() == k + !*/ + + void set_output_nr(long nr); + /*! + requires + - nr == -1 || nr > 0 + ensures + - #get_output_nr() == nr + !*/ + + void set_output_nc(long nc); + /*! + requires + - nc == -1 || nc > 0 + ensures + - #get_output_nc() == nc + !*/ + + template void setup(const SUBNET& sub); + /*! + requires + - SUBNET implements the SUBNET interface defined at the top of this file. + ensures + - Configures this layer to operate on the output of sub. + - If the total number of elements in the input tensor doesn't match the total + number of elements in the output tensor and the channel dimension is different, + an exception will be thrown. + !*/ + + template void forward(const SUBNET& sub, resizable_tensor& output); + /*! + requires + - SUBNET implements the SUBNET interface defined at the top of this file. + - setup() has been called. + ensures + - Reshapes or resizes the output of sub and stores it in #output. + - If is_spatial_rescale() == false, then performs a pure reshape operation. + - If is_spatial_rescale() == true, then performs bilinear interpolation to resize + the spatial dimensions while preserving the channel information. + - #output.num_samples() == sub.get_output().num_samples() + - #output.k() == get_output_k() if get_output_k() != -1, otherwise sub.get_output().k() + - #output.nr() == get_output_nr() if get_output_nr() != -1, otherwise sub.get_output().nr() + - #output.nc() == get_output_nc() if get_output_nc() != -1, otherwise sub.get_output().nc() + !*/ + + template void backward( + const tensor& gradient_input, + SUBNET& sub, + tensor& params_grad + ); + /*! + requires + - SUBNET implements the SUBNET interface defined at the top of this file. + - setup() has been called. + - gradient_input has the same dimensions as the output of forward(). + ensures + - Computes the gradients of this layer with respect to the input tensor and + parameters, and stores them in sub.get_gradient_input() and params_grad, + respectively. + - This function supports both pure reshaping and spatial rescaling operations. + !*/ + + dpoint map_input_to_output(dpoint p) const; + /*! + ensures + - Maps a point in the input tensor's coordinate system to the corresponding point + in the output tensor. This is useful for tracking how spatial locations change + through the network, especially during spatial rescaling. + !*/ + + dpoint map_output_to_input(dpoint p) const; + /*! + ensures + - Maps a point in the output tensor's coordinate system to the corresponding point + in the input tensor. This is the inverse of map_input_to_output(). + !*/ + + const tensor& get_layer_params() const; + /*! + ensures + - Returns the layer's parameters. This layer has no parameters, + so this always returns an empty tensor. + !*/ + + tensor& get_layer_params(); + /*! + ensures + - Returns the layer's parameters. This layer has no parameters, + so this always returns an empty tensor. + !*/ + }; + + template + using reshape_to = add_layer, SUBNET>; + + template + using flatten = add_layer, SUBNET>; + // ---------------------------------------------------------------------------------------- class dropout_ diff --git a/dlib/dnn/visitors.h b/dlib/dnn/visitors.h index 589e3556ef..d9f7401974 100644 --- a/dlib/dnn/visitors.h +++ b/dlib/dnn/visitors.h @@ -797,6 +797,15 @@ namespace dlib update(i); } + template + void operator()(size_t i, const add_layer, U, E>& l) + { + start_node(i, "linear"); + out << " | { outputs |{" << l.layer_details().get_num_outputs() << "}}"; + end_node(); + update(i); + } + template void operator()(size_t i, const add_layer&) { @@ -1031,6 +1040,20 @@ namespace dlib update(i); } + template + void operator()(size_t i, const add_layer, U, E>&) + { + start_node(i, "reshape_to"); + if (k == -1) out << " | {k|{unchanged}}"; + else out << " | {k|{" << k << "}}"; + if (nr == -1) out << " | {nr|{unchanged}}"; + else out << " | {nr|{" << nr << "}}"; + if (nc == -1) out << " | {nc|{unchanged}}"; + else out << " | {nc|{" << nc << "}}"; + end_node(); + update(i); + } + template void operator()(size_t i, const add_layer&) { diff --git a/dlib/test/dnn.cpp b/dlib/test/dnn.cpp index de26a1e1d3..35861a97c1 100644 --- a/dlib/test/dnn.cpp +++ b/dlib/test/dnn.cpp @@ -2538,7 +2538,19 @@ void test_embeddings() embeddings_<7, 12> l; auto res = test_layer(l); DLIB_TEST_MSG(res, res); - } + } + { + print_spinner(); + reshape_to_<-1, -1, -1> l; + auto res = test_layer(l); + DLIB_TEST_MSG(res, res); + } + { + print_spinner(); + reshape_to_<-1, 3, 5> l; + auto res = test_layer(l); + DLIB_TEST_MSG(res, res); + } } // ---------------------------------------------------------------------------------------- @@ -4801,6 +4813,39 @@ void test_multm_prev() } } + void test_resize_to() { + print_spinner(); + const long nr = 8, nc = 12; + const long n_samples = 5, k = 1, h = 4; + + using net_type = tag1>>>>>; + net_type net; + + dlib::rand rnd; + std::vector> x(n_samples); + matrix xtmp(nr, nc); + for (int ii = 0; ii < n_samples; ++ii) { + for (int jj = 0; jj < nr; ++jj) + for (int kk = 0; kk < nc; ++kk) + xtmp(jj, kk) = rnd.get_random_gaussian(); + x[ii] = xtmp; + } + + resizable_tensor input_tensor; + net.to_tensor(&x[0], &x[0] + n_samples, input_tensor); + net.forward(input_tensor); + + auto& output_tensor = layer(net).get_output(); + + DLIB_TEST(output_tensor.num_samples() == input_tensor.num_samples()); + DLIB_TEST(output_tensor.k() == input_tensor.k()); + DLIB_TEST(output_tensor.nr() == input_tensor.nr()); + DLIB_TEST(output_tensor.nc() == input_tensor.nc()); + DLIB_TEST(max(abs(mat(output_tensor) - mat(input_tensor))) < 1e-5); + } + // ---------------------------------------------------------------------------------------- template @@ -5170,6 +5215,7 @@ void test_multm_prev() test_embeddings(); test_tril(); test_basic_tensor_ops(); + test_resize_to(); test_layers(); test_visit_functions(); test_copy_tensor_cpu(); From 696586cc7271c20902f890fb7f08ff754f25e7f6 Mon Sep 17 00:00:00 2001 From: Davis King Date: Thu, 15 May 2025 07:40:46 -0400 Subject: [PATCH 21/39] tagging a new release --- dlib/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dlib/CMakeLists.txt b/dlib/CMakeLists.txt index dd67154f02..85a5e38d11 100644 --- a/dlib/CMakeLists.txt +++ b/dlib/CMakeLists.txt @@ -18,7 +18,7 @@ project(dlib) set(CPACK_PACKAGE_NAME "dlib") set(CPACK_PACKAGE_VERSION_MAJOR "19") set(CPACK_PACKAGE_VERSION_MINOR "24") -set(CPACK_PACKAGE_VERSION_PATCH "99") +set(CPACK_PACKAGE_VERSION_PATCH "9") set(VERSION ${CPACK_PACKAGE_VERSION_MAJOR}.${CPACK_PACKAGE_VERSION_MINOR}.${CPACK_PACKAGE_VERSION_PATCH}) # Only print these messages once, even if dlib is added multiple times via add_subdirectory() if (NOT TARGET dlib) From b9f5fa11ce90728a5bb1a29257d87c0c47e34d70 Mon Sep 17 00:00:00 2001 From: Davis King Date: Thu, 15 May 2025 07:41:16 -0400 Subject: [PATCH 22/39] put back to .99 --- dlib/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dlib/CMakeLists.txt b/dlib/CMakeLists.txt index 85a5e38d11..dd67154f02 100644 --- a/dlib/CMakeLists.txt +++ b/dlib/CMakeLists.txt @@ -18,7 +18,7 @@ project(dlib) set(CPACK_PACKAGE_NAME "dlib") set(CPACK_PACKAGE_VERSION_MAJOR "19") set(CPACK_PACKAGE_VERSION_MINOR "24") -set(CPACK_PACKAGE_VERSION_PATCH "9") +set(CPACK_PACKAGE_VERSION_PATCH "99") set(VERSION ${CPACK_PACKAGE_VERSION_MAJOR}.${CPACK_PACKAGE_VERSION_MINOR}.${CPACK_PACKAGE_VERSION_PATCH}) # Only print these messages once, even if dlib is added multiple times via add_subdirectory() if (NOT TARGET dlib) From 0a77590764ffebfa4268c4e9de8a5d25cd13fed0 Mon Sep 17 00:00:00 2001 From: Davis King Date: Sat, 24 May 2025 15:12:04 -0400 Subject: [PATCH 23/39] Make use_cuda() only set to true if there is also a GPU available --- dlib/cuda/cuda_dlib.cu | 7 + dlib/cuda/cuda_dlib.h | 6 + dlib/cuda/tensor_tools.cpp | 2 +- dlib/dnn/core_abstract.h | 8 +- dlib/test/dnn.cpp | 913 ++++++++++++++++++------------------- 5 files changed, 476 insertions(+), 460 deletions(-) diff --git a/dlib/cuda/cuda_dlib.cu b/dlib/cuda/cuda_dlib.cu index 5d6ec4052c..4aeb32bc58 100644 --- a/dlib/cuda/cuda_dlib.cu +++ b/dlib/cuda/cuda_dlib.cu @@ -44,6 +44,13 @@ namespace dlib CHECK_CUDA(cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync)); } + bool is_available( + ) + { + int num_devices; + return cudaGetDeviceCount(&num_devices) == cudaSuccess && num_devices > 0; + } + int get_num_devices ( ) { diff --git a/dlib/cuda/cuda_dlib.h b/dlib/cuda/cuda_dlib.h index 2f22b7e23e..2cc4c0c71e 100644 --- a/dlib/cuda/cuda_dlib.h +++ b/dlib/cuda/cuda_dlib.h @@ -25,6 +25,9 @@ namespace dlib int get_num_devices ( ); + bool is_available ( + ); + std::string get_device_name ( int device ); @@ -892,6 +895,9 @@ namespace dlib inline int get_num_devices ( ) { return 1; } + inline bool is_available ( + ) { return false; } + inline std::string get_device_name ( int device ) diff --git a/dlib/cuda/tensor_tools.cpp b/dlib/cuda/tensor_tools.cpp index b9818b3c73..c6640b30e0 100644 --- a/dlib/cuda/tensor_tools.cpp +++ b/dlib/cuda/tensor_tools.cpp @@ -22,7 +22,7 @@ namespace dlib ) { #ifdef DLIB_USE_CUDA - thread_local bool var(true); + thread_local bool var(cuda::is_available()); #else thread_local bool var(false); #endif diff --git a/dlib/dnn/core_abstract.h b/dlib/dnn/core_abstract.h index 70f1417e11..9d61c716b3 100644 --- a/dlib/dnn/core_abstract.h +++ b/dlib/dnn/core_abstract.h @@ -202,8 +202,12 @@ namespace dlib ensures - If dlib should use the CUDA implementation of a deep neural network then this function returns true and false otherwise. - - On program startup this function will return true if DLIB_USE_CUDA is defined. + - On program startup this function will return true if DLIB_USE_CUDA is defined and + there is an available GPU device to use. - This function always returns false if DLIB_USE_CUDA is not defined. + - This function sets a thread local variable. That is, each thread has its own value + for use_cuda(). This means that one thread may use cuda while another thread might + not, depending on the setting of use_cuda(). !*/ void set_use_cuda( @@ -211,7 +215,7 @@ namespace dlib ); /*! ensures - - #use_cuda() == flag + - #use_cuda() == flag for the calling thread. !*/ // ---------------------------------------------------------------------------------------- diff --git a/dlib/test/dnn.cpp b/dlib/test/dnn.cpp index 35861a97c1..6cd449b598 100644 --- a/dlib/test/dnn.cpp +++ b/dlib/test/dnn.cpp @@ -154,16 +154,16 @@ namespace dlog << LINFO << "src error: " << grad_error; DLIB_TEST(grad_error < 0.001); -#ifdef DLIB_USE_CUDA - resizable_tensor src1 = src; - resizable_tensor src2 = src; - resizable_tensor dest1, dest2; - dest1.copy_size(src); - dest2.copy_size(src); - cuda::softmax_all(dest1, src1); - cpu::softmax_all(dest2, src2); - DLIB_TEST_MSG(max(abs(mat(dest1)-mat(dest2))) < 1e-5, max(abs(mat(dest1)-mat(dest2)))); -#endif + IF_DLIB_USE_CUDA( + resizable_tensor src1 = src; + resizable_tensor src2 = src; + resizable_tensor dest1, dest2; + dest1.copy_size(src); + dest2.copy_size(src); + cuda::softmax_all(dest1, src1); + cpu::softmax_all(dest2, src2); + DLIB_TEST_MSG(max(abs(mat(dest1)-mat(dest2))) < 1e-5, max(abs(mat(dest1)-mat(dest2)))); + ) } void test_softmaxm() @@ -234,15 +234,16 @@ namespace cpu::softmax(output_tensor, input_tensor, operation_mode::PLANE_WISE); cpu::softmax_gradient(cpu_grad, output_tensor, gradient_input, operation_mode::PLANE_WISE); DLIB_TEST(max(abs(mat(output_tensor) - mat(expected_output))) < 1e-5); -#ifdef DLIB_USE_CUDA - resizable_tensor cuda_grad; - cuda_grad.copy_size(input_tensor); - cuda_grad = 0; - cuda::softmax(output_tensor, input_tensor, operation_mode::PLANE_WISE); - cpu::softmax_gradient(cuda_grad, output_tensor, gradient_input, operation_mode::PLANE_WISE); - DLIB_TEST(max(abs(mat(output_tensor) - mat(expected_output))) < 1e-5); - DLIB_TEST(max(abs(mat(cuda_grad) - mat(cpu_grad))) < 1e-5); -#endif + + IF_DLIB_USE_CUDA( + resizable_tensor cuda_grad; + cuda_grad.copy_size(input_tensor); + cuda_grad = 0; + cuda::softmax(output_tensor, input_tensor, operation_mode::PLANE_WISE); + cpu::softmax_gradient(cuda_grad, output_tensor, gradient_input, operation_mode::PLANE_WISE); + DLIB_TEST(max(abs(mat(output_tensor) - mat(expected_output))) < 1e-5); + DLIB_TEST(max(abs(mat(cuda_grad) - mat(cpu_grad))) < 1e-5); + ) } void test_softmax_all() @@ -284,222 +285,222 @@ namespace dlog << LINFO << "src error: " << grad_error; DLIB_TEST(grad_error < 0.001); -#ifdef DLIB_USE_CUDA - resizable_tensor src1 = src; - resizable_tensor src2 = src; - resizable_tensor dest1, dest2; - dest1.copy_size(src); - dest2.copy_size(src); - cuda::softmax_all(dest1, src1); - cpu::softmax_all(dest2, src2); - DLIB_TEST_MSG(max(abs(mat(dest1)-mat(dest2))) < 1e-5, max(abs(mat(dest1)-mat(dest2)))); -#endif + IF_DLIB_USE_CUDA( + resizable_tensor src1 = src; + resizable_tensor src2 = src; + resizable_tensor dest1, dest2; + dest1.copy_size(src); + dest2.copy_size(src); + cuda::softmax_all(dest1, src1); + cpu::softmax_all(dest2, src2); + DLIB_TEST_MSG(max(abs(mat(dest1)-mat(dest2))) < 1e-5, max(abs(mat(dest1)-mat(dest2)))); + ) } void test_mish() { -#ifdef DLIB_USE_CUDA - // make sure that cuda::mish and cpu::mish return the same results - using namespace dlib::tt; - print_spinner(); - const long n = 4; - const long k = 5; - const long nr = 3; - const long nc = 3; - resizable_tensor src(n,k,nr,nc); - tt::tensor_rand rnd; - rnd.fill_gaussian(src); + IF_DLIB_USE_CUDA( + // make sure that cuda::mish and cpu::mish return the same results + using namespace dlib::tt; + print_spinner(); + const long n = 4; + const long k = 5; + const long nr = 3; + const long nc = 3; + resizable_tensor src(n,k,nr,nc); + tt::tensor_rand rnd; + rnd.fill_gaussian(src); - resizable_tensor dest1, dest2; - dest1.copy_size(src); - dest2.copy_size(src); - // initialize to different values in order to make sure the output is actually changed - dest1 = 1; - dest2 = 2; - cuda::mish(dest1, src); - cpu::mish(dest2, src); - DLIB_TEST_MSG(max(abs(mat(dest1) - mat(dest2))) < 1e-6, max(abs(mat(dest1) - mat(dest2)))); -#endif // DLIB_USE_CUDA + resizable_tensor dest1, dest2; + dest1.copy_size(src); + dest2.copy_size(src); + // initialize to different values in order to make sure the output is actually changed + dest1 = 1; + dest2 = 2; + cuda::mish(dest1, src); + cpu::mish(dest2, src); + DLIB_TEST_MSG(max(abs(mat(dest1) - mat(dest2))) < 1e-6, max(abs(mat(dest1) - mat(dest2)))); + ) } void test_leaky_relu() { -#ifdef DLIB_USE_CUDA - using namespace dlib::tt; - print_spinner(); - const long n = 4; - const long k = 5; - const long nr = 3; - const long nc = 3; - const float alpha = 0.01; - resizable_tensor src(n, k, nr, nc); - tt::tensor_rand rnd; - rnd.fill_gaussian(src); - resizable_tensor dest_cuda, dest_cpu; - dest_cuda.copy_size(src); - dest_cpu.copy_size(src); - // initialize to different values in order to make sure the output is actually changed - dest_cuda = 1; - dest_cpu = 2; - cuda::leaky_relu(dest_cuda, src, alpha); - cpu::leaky_relu(dest_cpu, src, alpha); - - DLIB_TEST_MSG(max(abs(mat(dest_cuda) - mat(dest_cpu))) < 1e-7, max(abs(mat(dest_cuda) - mat(dest_cpu)))); -#endif // DLIB_USE_CUDA + IF_DLIB_USE_CUDA( + using namespace dlib::tt; + print_spinner(); + const long n = 4; + const long k = 5; + const long nr = 3; + const long nc = 3; + const float alpha = 0.01; + resizable_tensor src(n, k, nr, nc); + tt::tensor_rand rnd; + rnd.fill_gaussian(src); + resizable_tensor dest_cuda, dest_cpu; + dest_cuda.copy_size(src); + dest_cpu.copy_size(src); + // initialize to different values in order to make sure the output is actually changed + dest_cuda = 1; + dest_cpu = 2; + cuda::leaky_relu(dest_cuda, src, alpha); + cpu::leaky_relu(dest_cpu, src, alpha); + + DLIB_TEST_MSG(max(abs(mat(dest_cuda) - mat(dest_cpu))) < 1e-7, max(abs(mat(dest_cuda) - mat(dest_cpu)))); + ) } void test_clipped_relu() { -#ifdef DLIB_USE_CUDA - using namespace dlib::tt; - print_spinner(); - const long n = 4; - const long k = 5; - const long nr = 3; - const long nc = 3; - const float ceiling = 6.0f; - resizable_tensor src(n, k, nr, nc); - tt::tensor_rand rnd; - rnd.fill_gaussian(src, 0, 3); - resizable_tensor dest_cuda, dest_cpu; - dest_cuda.copy_size(src); - dest_cpu.copy_size(src); - // initialize to different values in order to make sure the output is actually changed - dest_cuda = 1; - dest_cpu = 2; - cuda::clipped_relu(dest_cuda, src, ceiling); - cpu::clipped_relu(dest_cpu, src, ceiling); - auto error = max(abs(mat(dest_cuda) - mat(dest_cpu))); - DLIB_TEST_MSG(error < 1e-7, "error: " << error); - - // test gradients - resizable_tensor grad_cuda, grad_cpu, grad_input; - grad_cuda.copy_size(src); - grad_cpu.copy_size(src); - grad_input.copy_size(src); - rnd.fill_uniform(grad_input); - grad_cuda = 0; - grad_cpu = 0; - cuda::clipped_relu_gradient(grad_cuda, dest_cuda, grad_input, ceiling); - cpu::clipped_relu_gradient(grad_cpu, dest_cpu, grad_input, ceiling); - error = max(abs(mat(grad_cuda) - mat(grad_cpu))); - DLIB_TEST_MSG(error < 1e-7, "error: " << error); -#endif // DLIB_USE_CUDA + IF_DLIB_USE_CUDA( + using namespace dlib::tt; + print_spinner(); + const long n = 4; + const long k = 5; + const long nr = 3; + const long nc = 3; + const float ceiling = 6.0f; + resizable_tensor src(n, k, nr, nc); + tt::tensor_rand rnd; + rnd.fill_gaussian(src, 0, 3); + resizable_tensor dest_cuda, dest_cpu; + dest_cuda.copy_size(src); + dest_cpu.copy_size(src); + // initialize to different values in order to make sure the output is actually changed + dest_cuda = 1; + dest_cpu = 2; + cuda::clipped_relu(dest_cuda, src, ceiling); + cpu::clipped_relu(dest_cpu, src, ceiling); + auto error = max(abs(mat(dest_cuda) - mat(dest_cpu))); + DLIB_TEST_MSG(error < 1e-7, "error: " << error); + + // test gradients + resizable_tensor grad_cuda, grad_cpu, grad_input; + grad_cuda.copy_size(src); + grad_cpu.copy_size(src); + grad_input.copy_size(src); + rnd.fill_uniform(grad_input); + grad_cuda = 0; + grad_cpu = 0; + cuda::clipped_relu_gradient(grad_cuda, dest_cuda, grad_input, ceiling); + cpu::clipped_relu_gradient(grad_cpu, dest_cpu, grad_input, ceiling); + error = max(abs(mat(grad_cuda) - mat(grad_cpu))); + DLIB_TEST_MSG(error < 1e-7, "error: " << error); + ) } void test_elu() { -#ifdef DLIB_USE_CUDA - using namespace dlib::tt; - print_spinner(); - const long n = 4; - const long k = 5; - const long nr = 3; - const long nc = 3; - const float alpha = 1.0f; - resizable_tensor src(n, k, nr, nc); - tt::tensor_rand rnd; - rnd.fill_gaussian(src); - resizable_tensor dest_cuda, dest_cpu; - dest_cuda.copy_size(src); - dest_cpu.copy_size(src); - // initialize to different values in order to make sure the output is actually changed - dest_cuda = 1; - dest_cpu = 2; - cuda::elu(dest_cuda, src, alpha); - cpu::elu(dest_cpu, src, alpha); - auto error = max(abs(mat(dest_cuda) - mat(dest_cpu))); - DLIB_TEST_MSG(error < 1e-7, "error: " << error); - // test gradients - resizable_tensor grad_cuda, grad_cpu, grad_input; - grad_cuda.copy_size(src); - grad_cpu.copy_size(src); - grad_input.copy_size(src); - rnd.fill_gaussian(grad_input); - grad_cuda = 0; - grad_cpu = 0; - cuda::elu_gradient(grad_cuda, dest_cuda, grad_input, alpha); - cpu::elu_gradient(grad_cpu, dest_cpu, grad_input, alpha); - error = max(abs(mat(grad_cuda) - mat(grad_cpu))); - DLIB_TEST_MSG(error < 1e-6, "error: " << error); -#endif // DLIB_USE_CUDA + IF_DLIB_USE_CUDA( + using namespace dlib::tt; + print_spinner(); + const long n = 4; + const long k = 5; + const long nr = 3; + const long nc = 3; + const float alpha = 1.0f; + resizable_tensor src(n, k, nr, nc); + tt::tensor_rand rnd; + rnd.fill_gaussian(src); + resizable_tensor dest_cuda, dest_cpu; + dest_cuda.copy_size(src); + dest_cpu.copy_size(src); + // initialize to different values in order to make sure the output is actually changed + dest_cuda = 1; + dest_cpu = 2; + cuda::elu(dest_cuda, src, alpha); + cpu::elu(dest_cpu, src, alpha); + auto error = max(abs(mat(dest_cuda) - mat(dest_cpu))); + DLIB_TEST_MSG(error < 1e-7, "error: " << error); + // test gradients + resizable_tensor grad_cuda, grad_cpu, grad_input; + grad_cuda.copy_size(src); + grad_cpu.copy_size(src); + grad_input.copy_size(src); + rnd.fill_gaussian(grad_input); + grad_cuda = 0; + grad_cpu = 0; + cuda::elu_gradient(grad_cuda, dest_cuda, grad_input, alpha); + cpu::elu_gradient(grad_cpu, dest_cpu, grad_input, alpha); + error = max(abs(mat(grad_cuda) - mat(grad_cpu))); + DLIB_TEST_MSG(error < 1e-6, "error: " << error); + ) } void test_gelu() { -#ifdef DLIB_USE_CUDA - // make sure that cuda::gelu and cpu::gelu return the same results - using namespace dlib::tt; - print_spinner(); - const long n = 4; - const long k = 5; - const long nr = 3; - const long nc = 3; - resizable_tensor src(n,k,nr,nc); - tt::tensor_rand rnd; - rnd.fill_gaussian(src); + IF_DLIB_USE_CUDA( + // make sure that cuda::gelu and cpu::gelu return the same results + using namespace dlib::tt; + print_spinner(); + const long n = 4; + const long k = 5; + const long nr = 3; + const long nc = 3; + resizable_tensor src(n,k,nr,nc); + tt::tensor_rand rnd; + rnd.fill_gaussian(src); - resizable_tensor dest1, dest2; - dest1.copy_size(src); - dest2.copy_size(src); - // initialize to different values in order to make sure the output is actually changed - dest1 = 1; - dest2 = 2; - cuda::gelu(dest1, src); - cpu::gelu(dest2, src); - DLIB_TEST_MSG(max(abs(mat(dest1) - mat(dest2))) < 1e-6, max(abs(mat(dest1) - mat(dest2)))); -#endif // DLIB_USE_CUDA + resizable_tensor dest1, dest2; + dest1.copy_size(src); + dest2.copy_size(src); + // initialize to different values in order to make sure the output is actually changed + dest1 = 1; + dest2 = 2; + cuda::gelu(dest1, src); + cpu::gelu(dest2, src); + DLIB_TEST_MSG(max(abs(mat(dest1) - mat(dest2))) < 1e-6, max(abs(mat(dest1) - mat(dest2)))); + ) } void test_smelu() { -#ifdef DLIB_USE_CUDA - using namespace dlib::tt; - print_spinner(); - const long n = 4; - const long k = 5; - const long nr = 3; - const long nc = 3; - const float beta = 1; - resizable_tensor src(n, k, nr, nc); - tt::tensor_rand rnd; - rnd.fill_gaussian(src); - resizable_tensor dest_cuda, dest_cpu; - dest_cuda.copy_size(src); - dest_cpu.copy_size(src); - // initialize to different values in order to make sure the output is actually changed - dest_cuda = 1; - dest_cpu = 2; - cuda::smelu(dest_cuda, src, beta); - cpu::smelu(dest_cpu, src, beta); - - DLIB_TEST_MSG(max(abs(mat(dest_cuda) - mat(dest_cpu))) < 1e-7, max(abs(mat(dest_cuda) - mat(dest_cpu)))); -#endif // DLIB_USE_CUDA + IF_DLIB_USE_CUDA( + using namespace dlib::tt; + print_spinner(); + const long n = 4; + const long k = 5; + const long nr = 3; + const long nc = 3; + const float beta = 1; + resizable_tensor src(n, k, nr, nc); + tt::tensor_rand rnd; + rnd.fill_gaussian(src); + resizable_tensor dest_cuda, dest_cpu; + dest_cuda.copy_size(src); + dest_cpu.copy_size(src); + // initialize to different values in order to make sure the output is actually changed + dest_cuda = 1; + dest_cpu = 2; + cuda::smelu(dest_cuda, src, beta); + cpu::smelu(dest_cpu, src, beta); + + DLIB_TEST_MSG(max(abs(mat(dest_cuda) - mat(dest_cpu))) < 1e-7, max(abs(mat(dest_cuda) - mat(dest_cpu)))); + ) } void test_silu() { -#ifdef DLIB_USE_CUDA - using namespace dlib::tt; - print_spinner(); - const long n = 4; - const long k = 5; - const long nr = 3; - const long nc = 3; - resizable_tensor src(n, k, nr, nc); - tt::tensor_rand rnd; - rnd.fill_gaussian(src); - resizable_tensor dest_cuda, dest_cpu; - dest_cuda.copy_size(src); - dest_cpu.copy_size(src); - // initialize to different values in order to make sure the output is actually changed - dest_cuda = 1; - dest_cpu = 2; - cuda::silu(dest_cuda, src); - cpu::silu(dest_cpu, src); - - DLIB_TEST_MSG(max(abs(mat(dest_cuda) - mat(dest_cpu))) < 1e-6, max(abs(mat(dest_cuda) - mat(dest_cpu)))); -#endif // DLIB_USE_CUDA + IF_DLIB_USE_CUDA( + using namespace dlib::tt; + print_spinner(); + const long n = 4; + const long k = 5; + const long nr = 3; + const long nc = 3; + resizable_tensor src(n, k, nr, nc); + tt::tensor_rand rnd; + rnd.fill_gaussian(src); + resizable_tensor dest_cuda, dest_cpu; + dest_cuda.copy_size(src); + dest_cpu.copy_size(src); + // initialize to different values in order to make sure the output is actually changed + dest_cuda = 1; + dest_cpu = 2; + cuda::silu(dest_cuda, src); + cpu::silu(dest_cpu, src); + + DLIB_TEST_MSG(max(abs(mat(dest_cuda) - mat(dest_cpu))) < 1e-6, max(abs(mat(dest_cuda) - mat(dest_cpu)))); + ) } void test_batch_normalize() @@ -710,28 +711,28 @@ namespace DLIB_TEST(::std::abs(rs.stddev() - 1.0f) < 0.01); } // check that the CPU and the CUDA implementation are equivalent -#ifdef DLIB_USE_CUDA - resizable_tensor y_cuda(x); - resizable_tensor means_cuda(x.num_samples()), invstds_cuda(x.num_samples()); - cuda::layer_normalize(eps, y_cuda, means_cuda, invstds_cuda, x, gamma, beta); - DLIB_TEST(max(abs(mat(y_cpu) - mat(y_cuda))) < 1e-5); - DLIB_TEST(max(abs(mat(means_cpu) - mat(means_cuda))) < 1e-5); - DLIB_TEST(max(abs(mat(invstds_cpu) - mat(invstds_cuda))) < 1e-5); - resizable_tensor gradient_input(x); - resizable_tensor src_grad_cpu(x), gamma_grad_cpu(1, x.k(), 1, 1), beta_grad_cpu(1, x.k(), 1, 1); - resizable_tensor src_grad_cuda(x), gamma_grad_cuda(1, x.k(), 1, 1), beta_grad_cuda(1, x.k(), 1, 1); - resizable_tensor dmeans_cpu, dvars_cpu, dmeans_cuda, dvars_cuda; - rnd.fill_gaussian(gradient_input); - src_grad_cpu = 0; - src_grad_cuda = 0; - cpu::layer_normalize_gradient(eps, gradient_input, means_cpu, invstds_cpu, x, gamma, src_grad_cpu, gamma_grad_cpu, beta_grad_cpu, dmeans_cpu, dvars_cpu); - cuda::layer_normalize_gradient(eps, gradient_input, means_cuda, invstds_cuda, x, gamma, src_grad_cuda, gamma_grad_cuda, beta_grad_cuda, dmeans_cuda, dvars_cuda); - DLIB_TEST(max(abs(mat(src_grad_cpu) - mat(src_grad_cuda))) < 1e-5); - DLIB_TEST(max(abs(mat(gamma_grad_cpu) - mat(gamma_grad_cuda))) < 1e-5); - DLIB_TEST(max(abs(mat(beta_grad_cpu) - mat(beta_grad_cuda))) < 1e-5); - DLIB_TEST(max(abs(mat(dmeans_cpu) - mat(dmeans_cuda))) < 1e-4); - DLIB_TEST(max(abs(mat(dvars_cpu) - mat(dvars_cuda))) < 1e-4); -#endif + IF_DLIB_USE_CUDA( + resizable_tensor y_cuda(x); + resizable_tensor means_cuda(x.num_samples()), invstds_cuda(x.num_samples()); + cuda::layer_normalize(eps, y_cuda, means_cuda, invstds_cuda, x, gamma, beta); + DLIB_TEST(max(abs(mat(y_cpu) - mat(y_cuda))) < 1e-5); + DLIB_TEST(max(abs(mat(means_cpu) - mat(means_cuda))) < 1e-5); + DLIB_TEST(max(abs(mat(invstds_cpu) - mat(invstds_cuda))) < 1e-5); + resizable_tensor gradient_input(x); + resizable_tensor src_grad_cpu(x), gamma_grad_cpu(1, x.k(), 1, 1), beta_grad_cpu(1, x.k(), 1, 1); + resizable_tensor src_grad_cuda(x), gamma_grad_cuda(1, x.k(), 1, 1), beta_grad_cuda(1, x.k(), 1, 1); + resizable_tensor dmeans_cpu, dvars_cpu, dmeans_cuda, dvars_cuda; + rnd.fill_gaussian(gradient_input); + src_grad_cpu = 0; + src_grad_cuda = 0; + cpu::layer_normalize_gradient(eps, gradient_input, means_cpu, invstds_cpu, x, gamma, src_grad_cpu, gamma_grad_cpu, beta_grad_cpu, dmeans_cpu, dvars_cpu); + cuda::layer_normalize_gradient(eps, gradient_input, means_cuda, invstds_cuda, x, gamma, src_grad_cuda, gamma_grad_cuda, beta_grad_cuda, dmeans_cuda, dvars_cuda); + DLIB_TEST(max(abs(mat(src_grad_cpu) - mat(src_grad_cuda))) < 1e-5); + DLIB_TEST(max(abs(mat(gamma_grad_cpu) - mat(gamma_grad_cuda))) < 1e-5); + DLIB_TEST(max(abs(mat(beta_grad_cpu) - mat(beta_grad_cuda))) < 1e-5); + DLIB_TEST(max(abs(mat(dmeans_cpu) - mat(dmeans_cuda))) < 1e-4); + DLIB_TEST(max(abs(mat(dvars_cpu) - mat(dvars_cuda))) < 1e-4); + ) } // ---------------------------------------------------------------------------------------- @@ -810,21 +811,21 @@ namespace DLIB_TEST(!backward_error_found); // check that the CPU and the CUDA implementation are equivalent -#ifdef DLIB_USE_CUDA - resizable_tensor y_cuda(x); - resizable_tensor scale_cuda; - cuda::rms_normalize(eps, y_cuda, scale_cuda, x, gamma); - DLIB_TEST(max(abs(mat(y_cpu) - mat(y_cuda))) < 1e-5); - DLIB_TEST(max(abs(mat(scale_cpu) - mat(scale_cuda))) < 1e-5); - - resizable_tensor src_grad_cuda(x), gamma_grad_cuda(1, x.k()); - resizable_tensor dscale_cuda(x.num_samples()); - src_grad_cuda = 0; - cuda::rms_normalize_gradient(gradient_input, scale_cuda, x, gamma, src_grad_cuda, gamma_grad_cuda, dscale_cuda); - DLIB_TEST(max(abs(mat(src_grad_cpu) - mat(src_grad_cuda))) < 1e-5); - DLIB_TEST(max(abs(mat(gamma_grad_cpu) - mat(gamma_grad_cuda))) < 1e-5); - DLIB_TEST(max(abs(mat(dscale_cpu) - mat(dscale_cuda))) < 1e-5); -#endif + IF_DLIB_USE_CUDA( + resizable_tensor y_cuda(x); + resizable_tensor scale_cuda; + cuda::rms_normalize(eps, y_cuda, scale_cuda, x, gamma); + DLIB_TEST(max(abs(mat(y_cpu) - mat(y_cuda))) < 1e-5); + DLIB_TEST(max(abs(mat(scale_cpu) - mat(scale_cuda))) < 1e-5); + + resizable_tensor src_grad_cuda(x), gamma_grad_cuda(1, x.k()); + resizable_tensor dscale_cuda(x.num_samples()); + src_grad_cuda = 0; + cuda::rms_normalize_gradient(gradient_input, scale_cuda, x, gamma, src_grad_cuda, gamma_grad_cuda, dscale_cuda); + DLIB_TEST(max(abs(mat(src_grad_cpu) - mat(src_grad_cuda))) < 1e-5); + DLIB_TEST(max(abs(mat(gamma_grad_cpu) - mat(gamma_grad_cuda))) < 1e-5); + DLIB_TEST(max(abs(mat(dscale_cpu) - mat(dscale_cuda))) < 1e-5); + ) } // ---------------------------------------------------------------------------------------- @@ -847,15 +848,15 @@ namespace input *= 2; DLIB_TEST(max(abs(mat(output_cpu_b) - mat(input))) < 1e-5); -#ifdef DLIB_USE_CUDA - input /= 2; - resizable_tensor output_cuda_a, output_cuda_b(input); - output_cuda_a.copy_size(output_cpu_a); - cuda::transpose(false, output_cuda_a, input); - cuda::transpose(true, output_cuda_b, output_cuda_a); - DLIB_TEST(max(abs(mat(output_cpu_a) - mat(output_cuda_a))) < 1e-5); - DLIB_TEST(max(abs(mat(output_cpu_b) - mat(output_cuda_b))) < 1e-5); -#endif + IF_DLIB_USE_CUDA( + input /= 2; + resizable_tensor output_cuda_a, output_cuda_b(input); + output_cuda_a.copy_size(output_cpu_a); + cuda::transpose(false, output_cuda_a, input); + cuda::transpose(true, output_cuda_b, output_cuda_a); + DLIB_TEST(max(abs(mat(output_cpu_a) - mat(output_cuda_a))) < 1e-5); + DLIB_TEST(max(abs(mat(output_cpu_b) - mat(output_cuda_b))) < 1e-5); + ) } // ---------------------------------------------------------------------------------------- @@ -1093,31 +1094,31 @@ void test_embeddings() memcpy(A, truth); DLIB_TEST(max(abs(mat(A)- mat(truth))) < 1e-5); -#ifdef DLIB_USE_CUDA - A = 4; - A.device(); - B.host(); - memcpy(A, truth); - DLIB_TEST(max(abs(mat(A)- mat(truth))) < 1e-5); + IF_DLIB_USE_CUDA( + A = 4; + A.device(); + B.host(); + memcpy(A, truth); + DLIB_TEST(max(abs(mat(A)- mat(truth))) < 1e-5); - A = 4; - A.device(); - B.device(); - memcpy(A, truth); - DLIB_TEST(max(abs(mat(A)- mat(truth))) < 1e-5); + A = 4; + A.device(); + B.device(); + memcpy(A, truth); + DLIB_TEST(max(abs(mat(A)- mat(truth))) < 1e-5); - A = 4; - A.host(); - B.device(); - memcpy(A, truth); - DLIB_TEST(max(abs(mat(A)- mat(truth))) < 1e-5); + A = 4; + A.host(); + B.device(); + memcpy(A, truth); + DLIB_TEST(max(abs(mat(A)- mat(truth))) < 1e-5); - A = 4; - A.host_write_only(); - B.device(); - memcpy(A, truth); - DLIB_TEST(max(abs(mat(A)- mat(truth))) < 1e-5); -#endif + A = 4; + A.host_write_only(); + B.device(); + memcpy(A, truth); + DLIB_TEST(max(abs(mat(A)- mat(truth))) < 1e-5); + ) } { @@ -1166,69 +1167,69 @@ void test_embeddings() } -#ifdef DLIB_USE_CUDA - A = 4; - A.device(); - B.host(); - { - // non-aliasing test - auto aA = at(A,5); - auto aB = at(B,5); - memcpy(aA, aB); - truth = {4,4,4,4,4, 1,1,1,1,1, 4}; - DLIB_TEST(max(abs(mat(A)- truth)) < 1e-5); - } - { - // aliasing test - auto aA = at(A,1); - auto aB = at(A,6); - memcpy(aA, aB); - truth = {4,1,1,1,1, 4,1,1,1,1, 4}; - DLIB_TEST(max(abs(mat(A)- truth)) < 1e-5); - } + IF_DLIB_USE_CUDA( + A = 4; + A.device(); + B.host(); + { + // non-aliasing test + auto aA = at(A,5); + auto aB = at(B,5); + memcpy(aA, aB); + truth = {4,4,4,4,4, 1,1,1,1,1, 4}; + DLIB_TEST(max(abs(mat(A)- truth)) < 1e-5); + } + { + // aliasing test + auto aA = at(A,1); + auto aB = at(A,6); + memcpy(aA, aB); + truth = {4,1,1,1,1, 4,1,1,1,1, 4}; + DLIB_TEST(max(abs(mat(A)- truth)) < 1e-5); + } - A = 4; - A.device(); - B.device(); - { - // non-aliasing test - auto aA = at(A,5); - auto aB = at(B,5); - memcpy(aA, aB); - truth = {4,4,4,4,4, 1,1,1,1,1, 4}; - DLIB_TEST(max(abs(mat(A)- truth)) < 1e-5); - } - { - // aliasing test - auto aA = at(A,1); - auto aB = at(A,6); - memcpy(aA, aB); - truth = {4,1,1,1,1, 4,1,1,1,1, 4}; - DLIB_TEST(max(abs(mat(A)- truth)) < 1e-5); - } + A = 4; + A.device(); + B.device(); + { + // non-aliasing test + auto aA = at(A,5); + auto aB = at(B,5); + memcpy(aA, aB); + truth = {4,4,4,4,4, 1,1,1,1,1, 4}; + DLIB_TEST(max(abs(mat(A)- truth)) < 1e-5); + } + { + // aliasing test + auto aA = at(A,1); + auto aB = at(A,6); + memcpy(aA, aB); + truth = {4,1,1,1,1, 4,1,1,1,1, 4}; + DLIB_TEST(max(abs(mat(A)- truth)) < 1e-5); + } - A = 4; - A.host(); - B.device(); - { - // non-aliasing test - auto aA = at(A,5); - auto aB = at(B,5); - memcpy(aA, aB); - truth = {4,4,4,4,4, 1,1,1,1,1, 4}; - DLIB_TEST(max(abs(mat(A)- truth)) < 1e-5); - } - { - // aliasing test - auto aA = at(A,1); - auto aB = at(A,6); - memcpy(aA, aB); - truth = {4,1,1,1,1, 4,1,1,1,1, 4}; - DLIB_TEST(max(abs(mat(A)- truth)) < 1e-5); - } + A = 4; + A.host(); + B.device(); + { + // non-aliasing test + auto aA = at(A,5); + auto aB = at(B,5); + memcpy(aA, aB); + truth = {4,4,4,4,4, 1,1,1,1,1, 4}; + DLIB_TEST(max(abs(mat(A)- truth)) < 1e-5); + } + { + // aliasing test + auto aA = at(A,1); + auto aB = at(A,6); + memcpy(aA, aB); + truth = {4,1,1,1,1, 4,1,1,1,1, 4}; + DLIB_TEST(max(abs(mat(A)- truth)) < 1e-5); + } -#endif + ) } { @@ -3672,22 +3673,22 @@ void test_multm_prev() trainer.train(inputs, labels); const auto error_after = compute_error(); DLIB_TEST_MSG(error_after < error_before, "multi channel error increased after training"); -#if DLIB_USE_CUDA - cuda::compute_loss_mean_squared_per_channel_and_pixel cuda_compute; - cpu::compute_loss_mean_squared_per_channel_and_pixel cpu_compute; - double cuda_loss, cpu_loss; - const tensor& output_tensor = net.subnet().get_output(); - resizable_tensor cuda_grad(output_tensor), cpu_grad(output_tensor); - cuda_compute(labels.begin(), output_tensor, cuda_grad, cuda_loss); - cpu_compute(labels.begin(), output_tensor, cpu_grad, cpu_loss); - DLIB_TEST(cuda_grad.size() == cpu_grad.size()); - for (size_t i = 0; i < cuda_grad.size(); ++i) - { - DLIB_TEST(::std::abs(*(cuda_grad.begin() + i) - *(cpu_grad.begin() + i)) < 1e-8); - } - const auto err = abs(cuda_loss - cpu_loss) / cpu_loss; - DLIB_TEST_MSG(err < 1e-6, "multi channel cuda and cpu losses differ"); -#endif + IF_DLIB_USE_CUDA( + cuda::compute_loss_mean_squared_per_channel_and_pixel cuda_compute; + cpu::compute_loss_mean_squared_per_channel_and_pixel cpu_compute; + double cuda_loss, cpu_loss; + const tensor& output_tensor = net.subnet().get_output(); + resizable_tensor cuda_grad(output_tensor), cpu_grad(output_tensor); + cuda_compute(labels.begin(), output_tensor, cuda_grad, cuda_loss); + cpu_compute(labels.begin(), output_tensor, cpu_grad, cpu_loss); + DLIB_TEST(cuda_grad.size() == cpu_grad.size()); + for (size_t i = 0; i < cuda_grad.size(); ++i) + { + DLIB_TEST(::std::abs(*(cuda_grad.begin() + i) - *(cpu_grad.begin() + i)) < 1e-8); + } + const auto err = abs(cuda_loss - cpu_loss) / cpu_loss; + DLIB_TEST_MSG(err < 1e-6, "multi channel cuda and cpu losses differ"); + ) } // ---------------------------------------------------------------------------------------- @@ -3883,22 +3884,22 @@ void test_multm_prev() DLIB_TEST_MSG(num_correct >= num_correct_required, "Number of correctly classified elements = " << num_correct << ", required = " << num_correct_required); -#if DLIB_USE_CUDA - cuda::compute_loss_binary_log_per_pixel cuda_compute; - cpu::compute_loss_binary_log_per_pixel cpu_compute; - double cuda_loss, cpu_loss; - const tensor& output_tensor = net.subnet().get_output(); - resizable_tensor cuda_grad(output_tensor), cpu_grad(output_tensor); - cuda_compute(y.begin(), output_tensor, cuda_grad, cuda_loss); - cpu_compute(y.begin(), output_tensor, cpu_grad, cpu_loss); - DLIB_TEST(cuda_grad.size() == cpu_grad.size()); - for (size_t i = 0; i < cuda_grad.size(); ++i) - { - DLIB_TEST(::std::abs(*(cuda_grad.begin() + i) - *(cpu_grad.begin() + i)) < 1e-8); - } - const auto err = abs(cuda_loss - cpu_loss) / cpu_loss; - DLIB_TEST_MSG(err < 1e-6, "binary log per pixel cuda and cpu losses differ"); -#endif + IF_DLIB_USE_CUDA( + cuda::compute_loss_binary_log_per_pixel cuda_compute; + cpu::compute_loss_binary_log_per_pixel cpu_compute; + double cuda_loss, cpu_loss; + const tensor& output_tensor = net.subnet().get_output(); + resizable_tensor cuda_grad(output_tensor), cpu_grad(output_tensor); + cuda_compute(y.begin(), output_tensor, cuda_grad, cuda_loss); + cpu_compute(y.begin(), output_tensor, cpu_grad, cpu_loss); + DLIB_TEST(cuda_grad.size() == cpu_grad.size()); + for (size_t i = 0; i < cuda_grad.size(); ++i) + { + DLIB_TEST(::std::abs(*(cuda_grad.begin() + i) - *(cpu_grad.begin() + i)) < 1e-8); + } + const auto err = abs(cuda_loss - cpu_loss) / cpu_loss; + DLIB_TEST_MSG(err < 1e-6, "binary log per pixel cuda and cpu losses differ"); + ) } // ---------------------------------------------------------------------------------------- @@ -4234,22 +4235,22 @@ void test_multm_prev() DLIB_TEST_MSG(num_correct >= num_correct_required, "Number of correctly classified elements = " << num_correct << ", required = " << num_correct_required); -#if DLIB_USE_CUDA - cuda::compute_loss_multiclass_log_per_pixel cuda_compute; - cpu::compute_loss_multiclass_log_per_pixel cpu_compute; - double cuda_loss, cpu_loss; - const tensor& output_tensor = net.subnet().get_output(); - resizable_tensor cuda_grad(output_tensor), cpu_grad(output_tensor); - cuda_compute(y.begin(), output_tensor, cuda_grad, cuda_loss); - cpu_compute(y.begin(), output_tensor, cpu_grad, cpu_loss); - DLIB_TEST(cuda_grad.size() == cpu_grad.size()); - for (size_t i = 0; i < cuda_grad.size(); ++i) - { - DLIB_TEST(::std::abs(*(cuda_grad.begin() + i) - *(cpu_grad.begin() + i)) < 1e-8); - } - const auto err = abs(cuda_loss - cpu_loss) / cpu_loss; - DLIB_TEST_MSG(err < 1e-6, "multiclass log per pixel cuda and cpu losses differ"); -#endif + IF_DLIB_USE_CUDA( + cuda::compute_loss_multiclass_log_per_pixel cuda_compute; + cpu::compute_loss_multiclass_log_per_pixel cpu_compute; + double cuda_loss, cpu_loss; + const tensor& output_tensor = net.subnet().get_output(); + resizable_tensor cuda_grad(output_tensor), cpu_grad(output_tensor); + cuda_compute(y.begin(), output_tensor, cuda_grad, cuda_loss); + cpu_compute(y.begin(), output_tensor, cpu_grad, cpu_loss); + DLIB_TEST(cuda_grad.size() == cpu_grad.size()); + for (size_t i = 0; i < cuda_grad.size(); ++i) + { + DLIB_TEST(::std::abs(*(cuda_grad.begin() + i) - *(cpu_grad.begin() + i)) < 1e-8); + } + const auto err = abs(cuda_loss - cpu_loss) / cpu_loss; + DLIB_TEST_MSG(err < 1e-6, "multiclass log per pixel cuda and cpu losses differ"); + ) } // ---------------------------------------------------------------------------------------- @@ -4345,22 +4346,22 @@ void test_multm_prev() "The weighted class (" << weighted_class << ") does not dominate: " << num_weighted_class << " <= " << num_not_weighted_class); -#if DLIB_USE_CUDA - cuda::compute_loss_multiclass_log_per_pixel_weighted cuda_compute; - cpu::compute_loss_multiclass_log_per_pixel_weighted cpu_compute; - double cuda_loss, cpu_loss; - const tensor& output_tensor = net.subnet().get_output(); - resizable_tensor cuda_grad(output_tensor), cpu_grad(output_tensor); - cuda_compute(y_weighted.begin(), output_tensor, cuda_grad, cuda_loss); - cpu_compute(y_weighted.begin(), output_tensor, cpu_grad, cpu_loss); - DLIB_TEST(cuda_grad.size() == cpu_grad.size()); - for (size_t i = 0; i < cuda_grad.size(); ++i) - { - DLIB_TEST(::std::abs(*(cuda_grad.begin() + i) - *(cpu_grad.begin() + i)) < 1e-8); - } - const auto err = abs(cuda_loss - cpu_loss) / cpu_loss; - DLIB_TEST_MSG(err < 1e-5, "multi class log per pixel weighted cuda and cpu losses differ: " << err); -#endif + IF_DLIB_USE_CUDA( + cuda::compute_loss_multiclass_log_per_pixel_weighted cuda_compute; + cpu::compute_loss_multiclass_log_per_pixel_weighted cpu_compute; + double cuda_loss, cpu_loss; + const tensor& output_tensor = net.subnet().get_output(); + resizable_tensor cuda_grad(output_tensor), cpu_grad(output_tensor); + cuda_compute(y_weighted.begin(), output_tensor, cuda_grad, cuda_loss); + cpu_compute(y_weighted.begin(), output_tensor, cpu_grad, cpu_loss); + DLIB_TEST(cuda_grad.size() == cpu_grad.size()); + for (size_t i = 0; i < cuda_grad.size(); ++i) + { + DLIB_TEST(::std::abs(*(cuda_grad.begin() + i) - *(cpu_grad.begin() + i)) < 1e-8); + } + const auto err = abs(cuda_loss - cpu_loss) / cpu_loss; + DLIB_TEST_MSG(err < 1e-5, "multi class log per pixel weighted cuda and cpu losses differ: " << err); + ) } } @@ -4527,10 +4528,10 @@ void test_multm_prev() img = 1; img.host()[idx] = 2; cpu::resize_bilinear(out, img); -#ifdef DLIB_USE_CUDA - cuda::resize_bilinear(out2, img); - DLIB_TEST(max(abs(mat(out)-mat(out2))) < 1e-5); -#endif + IF_DLIB_USE_CUDA( + cuda::resize_bilinear(out2, img); + DLIB_TEST(max(abs(mat(out)-mat(out2))) < 1e-5); + ) resizable_tensor gradient_input; gradient_input.copy_size(out); @@ -4561,12 +4562,12 @@ void test_multm_prev() dlog << LINFO << "analytic grad: "<< grad2.host()[idx]-0.1; DLIB_TEST_MSG(std::abs(numerical_grad - grad2.host()[idx]+0.1) < 1e-2, std::abs(numerical_grad - grad2.host()[idx]+0.1) << " numerical_grad: " << numerical_grad); -#ifdef DLIB_USE_CUDA - cuda::resize_bilinear_gradient(grad, gradient_input); - dlog << LINFO << "analytic grad: "<< grad.host()[idx]-0.1; - DLIB_TEST_MSG(std::abs(numerical_grad - grad.host()[idx]+0.1) < 1e-2, std::abs(numerical_grad - grad.host()[idx]+0.1) << " numerical_grad: " << numerical_grad); - DLIB_TEST(max(abs(mat(grad)-mat(grad2))) < 1e-5); -#endif + IF_DLIB_USE_CUDA( + cuda::resize_bilinear_gradient(grad, gradient_input); + dlog << LINFO << "analytic grad: "<< grad.host()[idx]-0.1; + DLIB_TEST_MSG(std::abs(numerical_grad - grad.host()[idx]+0.1) < 1e-2, std::abs(numerical_grad - grad.host()[idx]+0.1) << " numerical_grad: " << numerical_grad); + DLIB_TEST(max(abs(mat(grad)-mat(grad2))) < 1e-5); + ) } @@ -4587,11 +4588,11 @@ void test_multm_prev() auto wout = aout(out, out.nc()*1+1); auto wimg = aimg(img, img.nc()*1+1); cpu::resize_bilinear(wout,out.nc(),out.nr()*out.nc(), wimg,img.nc(),img.nr()*img.nc()); -#ifdef DLIB_USE_CUDA - auto wout2 = aout(out2, out2.nc()*1+1); - cuda::resize_bilinear(wout2,out2.nc(),out2.nr()*out2.nc(), wimg,img.nc(),img.nr()*img.nc()); - DLIB_TEST(max(abs(mat(out)-mat(out2))) < 1e-5); -#endif + IF_DLIB_USE_CUDA( + auto wout2 = aout(out2, out2.nc()*1+1); + cuda::resize_bilinear(wout2,out2.nc(),out2.nr()*out2.nc(), wimg,img.nc(),img.nr()*img.nc()); + DLIB_TEST(max(abs(mat(out)-mat(out2))) < 1e-5); + ) resizable_tensor gradient_input; @@ -4629,15 +4630,14 @@ void test_multm_prev() dlog << LINFO << "analytic grad: "<< grad2.host()[idx]-0.1; DLIB_TEST_MSG(std::abs(numerical_grad - grad2.host()[idx]+0.1) < 1e-2, std::abs(numerical_grad - grad2.host()[idx]+0.1) << " numerical_grad: " << numerical_grad); -#ifdef DLIB_USE_CUDA - wgrad2 = aimg(grad, grad.nc()*1+1); - wgradient_input = aout(gradient_input, gradient_input.nc()*1+1); - cuda::resize_bilinear_gradient(wgrad2,grad.nc(),grad.nr()*grad.nc(), wgradient_input,gradient_input.nc(),gradient_input.nr()*gradient_input.nc()); - dlog << LINFO << "analytic grad: "<< grad.host()[idx]-0.1; - DLIB_TEST_MSG(std::abs(numerical_grad - grad.host()[idx]+0.1) < 1e-2, std::abs(numerical_grad - grad.host()[idx]+0.1) << " numerical_grad: " << numerical_grad); - DLIB_TEST_MSG(max(abs(mat(grad)-mat(grad2))) < 1e-5, max(abs(mat(grad)-mat(grad2)))); -#endif - + IF_DLIB_USE_CUDA( + wgrad2 = aimg(grad, grad.nc()*1+1); + wgradient_input = aout(gradient_input, gradient_input.nc()*1+1); + cuda::resize_bilinear_gradient(wgrad2,grad.nc(),grad.nr()*grad.nc(), wgradient_input,gradient_input.nc(),gradient_input.nr()*gradient_input.nc()); + dlog << LINFO << "analytic grad: "<< grad.host()[idx]-0.1; + DLIB_TEST_MSG(std::abs(numerical_grad - grad.host()[idx]+0.1) < 1e-2, std::abs(numerical_grad - grad.host()[idx]+0.1) << " numerical_grad: " << numerical_grad); + DLIB_TEST_MSG(max(abs(mat(grad)-mat(grad2))) < 1e-5, max(abs(mat(grad)-mat(grad2)))); + ) } } @@ -5037,20 +5037,20 @@ void test_multm_prev() void test_reorg() { -#ifdef DLIB_USE_CUDA - print_spinner(); - resizable_tensor x(2, 4, 8, 16); - resizable_tensor out_cpu(2, 16, 4, 8), out_cuda(2, 16, 4, 8); - resizable_tensor grad_cpu(x), grad_cuda(x); - tt::tensor_rand rnd; - rnd.fill_gaussian(x); - cpu::reorg(false, out_cpu, 2, 2, x); - cuda::reorg(false, out_cuda, 2, 2, x); - DLIB_TEST(max(squared(mat(out_cuda) - mat(out_cpu))) == 0); - cpu::reorg_gradient(false, grad_cpu, 2, 2, out_cpu); - cuda::reorg_gradient(false, grad_cuda, 2, 2, out_cuda); - DLIB_TEST(max(squared(mat(out_cuda) - mat(out_cpu))) == 0); -#endif + IF_DLIB_USE_CUDA( + print_spinner(); + resizable_tensor x(2, 4, 8, 16); + resizable_tensor out_cpu(2, 16, 4, 8), out_cuda(2, 16, 4, 8); + resizable_tensor grad_cpu(x), grad_cuda(x); + tt::tensor_rand rnd; + rnd.fill_gaussian(x); + cpu::reorg(false, out_cpu, 2, 2, x); + cuda::reorg(false, out_cuda, 2, 2, x); + DLIB_TEST(max(squared(mat(out_cuda) - mat(out_cpu))) == 0); + cpu::reorg_gradient(false, grad_cpu, 2, 2, out_cpu); + cuda::reorg_gradient(false, grad_cuda, 2, 2, out_cuda); + DLIB_TEST(max(squared(mat(out_cuda) - mat(out_cpu))) == 0); + ) } void test_input_tensor() @@ -5144,29 +5144,28 @@ void test_multm_prev() srand(1234); test_tagging(); -#ifdef DLIB_USE_CUDA - set_use_cuda(true); - test_affine_rect(); - test_conv(); - test_more_ops2(); - test_more_ops(1,1); - test_more_ops(3,4); - test_more_ops(4,3); - test_more_ops(4,1); - test_more_ops(1,4); - test_more_ops(10000,4); - compare_bn_gpu_and_cpu(); - compare_bn_conv_gpu_and_cpu(); - test_add(); - test_multiply_zero_padded(); - compare_adam(); - test_copy_tensor_gpu(); - test_copy_tensor_add_to_gpu(); - test_copy_tensor_gpu(); - test_copy_tensor_add_to_gpu(); - test_scale_channels(); -#endif - set_use_cuda(false); + IF_DLIB_USE_CUDA( + test_affine_rect(); + test_conv(); + test_more_ops2(); + test_more_ops(1,1); + test_more_ops(3,4); + test_more_ops(4,3); + test_more_ops(4,1); + test_more_ops(1,4); + test_more_ops(10000,4); + compare_bn_gpu_and_cpu(); + compare_bn_conv_gpu_and_cpu(); + test_add(); + test_multiply_zero_padded(); + compare_adam(); + test_copy_tensor_gpu(); + test_copy_tensor_add_to_gpu(); + test_copy_tensor_gpu(); + test_copy_tensor_add_to_gpu(); + test_scale_channels(); + ) + test_tensor_resize_bilinear(2, 3, 6,6, 11, 11); test_tensor_resize_bilinear(2, 3, 6,6, 3, 4); test_tensor_resize_bilinear(2, 3, 5,6, 12, 21); From 345b9b7c5e5182b2eb6c79a025508ca6ce08030b Mon Sep 17 00:00:00 2001 From: Kane Scipioni Date: Fri, 18 Jul 2025 00:08:51 -0500 Subject: [PATCH 24/39] Allocate CUDA memory only when use_cuda() returns true --- dlib/cuda/cuda_dlib.cu | 33 ++++++++++++++++++-- dlib/cuda/cuda_dlib.h | 14 +++++++++ dlib/cuda/curand_dlibapi.cpp | 12 +++++--- dlib/cuda/gpu_data.cpp | 22 +++++++++++-- dlib/cuda/gpu_data.h | 19 ++++++++++-- dlib/cuda/tensor.h | 50 +++++++++++++++++++----------- dlib/cuda/tensor_tools.cpp | 26 ---------------- dlib/cuda/tensor_tools.h | 15 --------- dlib/dnn/loss.h | 60 +++++++++++++++++++----------------- 9 files changed, 151 insertions(+), 100 deletions(-) diff --git a/dlib/cuda/cuda_dlib.cu b/dlib/cuda/cuda_dlib.cu index 4aeb32bc58..921bbdbd35 100644 --- a/dlib/cuda/cuda_dlib.cu +++ b/dlib/cuda/cuda_dlib.cu @@ -9,6 +9,17 @@ namespace dlib { + namespace + { + bool& use_cuda_impl ( + ) + { + thread_local bool var(cuda::is_available()); + return var; + } + + } + namespace cuda { @@ -18,14 +29,16 @@ namespace dlib int dev ) { - CHECK_CUDA(cudaSetDevice(dev)); + if (is_available()) + CHECK_CUDA(cudaSetDevice(dev)); } int get_device ( ) { - int dev = 0; - CHECK_CUDA(cudaGetDevice(&dev)); + int dev = -1; + if (is_available()) + CHECK_CUDA(cudaGetDevice(&dev)); return dev; } @@ -51,6 +64,20 @@ namespace dlib return cudaGetDeviceCount(&num_devices) == cudaSuccess && num_devices > 0; } + bool use_cuda( + ) + { + return use_cuda_impl(); + } + + void set_use_cuda( + bool flag + ) + { + if (is_available()) + use_cuda_impl() = flag; + } + int get_num_devices ( ) { diff --git a/dlib/cuda/cuda_dlib.h b/dlib/cuda/cuda_dlib.h index 2cc4c0c71e..289b2c6ec3 100644 --- a/dlib/cuda/cuda_dlib.h +++ b/dlib/cuda/cuda_dlib.h @@ -28,6 +28,13 @@ namespace dlib bool is_available ( ); + bool use_cuda( + ); + + void set_use_cuda( + bool flag + ); + std::string get_device_name ( int device ); @@ -898,6 +905,13 @@ namespace dlib inline bool is_available ( ) { return false; } + inline bool use_cuda( + ) { return false; } + + inline void set_use_cuda( + bool flag + ) {} + inline std::string get_device_name ( int device ) diff --git a/dlib/cuda/curand_dlibapi.cpp b/dlib/cuda/curand_dlibapi.cpp index 67828e6640..6f8b00006b 100644 --- a/dlib/cuda/curand_dlibapi.cpp +++ b/dlib/cuda/curand_dlibapi.cpp @@ -6,6 +6,7 @@ #ifdef DLIB_USE_CUDA #include "curand_dlibapi.h" +#include "cuda_dlib.h" #include #include "../string.h" @@ -47,11 +48,14 @@ namespace dlib unsigned long long seed ) : handle(nullptr) { - curandGenerator_t gen; - CHECK_CURAND(curandCreateGenerator(&gen, CURAND_RNG_PSEUDO_DEFAULT)); - handle = gen; + if (is_available()) + { + curandGenerator_t gen; + CHECK_CURAND(curandCreateGenerator(&gen, CURAND_RNG_PSEUDO_DEFAULT)); + handle = gen; - CHECK_CURAND(curandSetPseudoRandomGeneratorSeed(gen, seed)); + CHECK_CURAND(curandSetPseudoRandomGeneratorSeed(gen, seed)); + } } curand_generator:: diff --git a/dlib/cuda/gpu_data.cpp b/dlib/cuda/gpu_data.cpp index 64f184aede..413c3996ee 100644 --- a/dlib/cuda/gpu_data.cpp +++ b/dlib/cuda/gpu_data.cpp @@ -58,10 +58,16 @@ namespace dlib if (dest_offset == 0 && num == dest.size()) { // copy the memory efficiently based on which copy is current in each object. - if (src.device_ready()) + if (dest.device_id() >= 0 && src.device_ready()) CHECK_CUDA(cudaMemcpy(dest.device_write_only(), src.device()+src_offset, num*sizeof(float), cudaMemcpyDeviceToDevice)); - else + else if (dest.device_id() < 0 && src.device_ready()) + CHECK_CUDA(cudaMemcpy(dest.host_write_only(), src.device()+src_offset, num*sizeof(float), cudaMemcpyDeviceToHost)); + else if (dest.device_id() >= 0 && !src.device_ready()) CHECK_CUDA(cudaMemcpy(dest.device_write_only(), src.host()+src_offset, num*sizeof(float), cudaMemcpyHostToDevice)); + else if (dest.device_id() >= 0 || src.device_id() >= 0) + CHECK_CUDA(cudaMemcpy(dest.host_write_only(), src.host()+src_offset, num*sizeof(float), cudaMemcpyHostToHost)); + else + std::memcpy(dest.host_write_only(), src.host()+src_offset, num*sizeof(float)); } else { @@ -72,8 +78,11 @@ namespace dlib CHECK_CUDA(cudaMemcpy(dest.host()+dest_offset, src.device()+src_offset, num*sizeof(float), cudaMemcpyDeviceToHost)); else if (dest.device_ready() && !src.device_ready()) CHECK_CUDA(cudaMemcpy(dest.device()+dest_offset, src.host()+src_offset, num*sizeof(float), cudaMemcpyHostToDevice)); - else + else if (dest.device_id() >= 0 || src.device_id() >= 0) CHECK_CUDA(cudaMemcpy(dest.host()+dest_offset, src.host()+src_offset, num*sizeof(float), cudaMemcpyHostToHost)); + else + std::memcpy(dest.host()+dest_offset, src.host()+src_offset, num*sizeof(float)); + } } } @@ -199,6 +208,13 @@ namespace dlib device_current = true; device_in_use = false; + if (!cuda::use_cuda()) + { + data_host.reset(new float[new_size], std::default_delete()); + the_device_id = -1; + return; + } + try { CHECK_CUDA(cudaGetDevice(&the_device_id)); diff --git a/dlib/cuda/gpu_data.h b/dlib/cuda/gpu_data.h index 022a05f71c..4d37315a40 100644 --- a/dlib/cuda/gpu_data.h +++ b/dlib/cuda/gpu_data.h @@ -12,6 +12,14 @@ namespace dlib { +// ---------------------------------------------------------------------------------------- + + namespace cuda + { + bool use_cuda( + ); + } + // ---------------------------------------------------------------------------------------- class gpu_data @@ -93,14 +101,16 @@ namespace dlib float* host() { copy_to_host(); - device_current = false; + if (device_id() >= 0) + device_current = false; return data_host.get(); } float* host_write_only() { host_current = true; - device_current = false; + if (device_id() >= 0) + device_current = false; return data_host.get(); } @@ -109,6 +119,7 @@ namespace dlib #ifndef DLIB_USE_CUDA DLIB_CASSERT(false, "CUDA NOT ENABLED"); #endif + DLIB_CASSERT(device_id() >= 0, "This data is host only"); copy_to_device(); device_in_use = true; return data_device.get(); @@ -119,6 +130,7 @@ namespace dlib #ifndef DLIB_USE_CUDA DLIB_CASSERT(false, "CUDA NOT ENABLED"); #endif + DLIB_CASSERT(device_id() >= 0, "This data is host only"); copy_to_device(); host_current = false; device_in_use = true; @@ -130,6 +142,7 @@ namespace dlib #ifndef DLIB_USE_CUDA DLIB_CASSERT(false, "CUDA NOT ENABLED"); #endif + DLIB_CASSERT(device_id() >= 0, "This data is host only"); wait_for_transfer_to_finish(); host_current = false; device_current = true; @@ -141,7 +154,7 @@ namespace dlib ) const { return host_current; } bool device_ready ( - ) const { return device_current && !have_active_transfer; } + ) const { return device_current && !have_active_transfer && device_id() >= 0; } size_t size() const { return data_size; } diff --git a/dlib/cuda/tensor.h b/dlib/cuda/tensor.h index 6a893df311..aaae4e836e 100644 --- a/dlib/cuda/tensor.h +++ b/dlib/cuda/tensor.h @@ -12,6 +12,18 @@ #include #include "../any.h" +#ifdef DLIB_USE_CUDA +#define IF_DLIB_USE_CUDA(...) if (cuda::use_cuda()) { __VA_ARGS__ } +#else +#define IF_DLIB_USE_CUDA(...) +#endif + +#ifdef DLIB_USE_CUDA +#define IF_DLIB_NOT_USE_CUDA(...) if (!cuda::use_cuda()) { __VA_ARGS__ } +#else +#define IF_DLIB_NOT_USE_CUDA(...) __VA_ARGS__ +#endif + namespace dlib { @@ -77,17 +89,18 @@ namespace dlib tensor& operator= (float val) { -#ifdef DLIB_USE_CUDA - // If you are using CUDA then presumably you will be mostly using tensors on - // the GPU. So unless you seem to be actively working with the host side's - // data then we do this initialization on the device side since this avoids a - // host to device transfer that would likely immediately follow. - if (data().device_ready()) - { - cuda::set_tensor(*this, val); - return *this; - } -#endif + IF_DLIB_USE_CUDA( + // If you are using CUDA then presumably you will be mostly using tensors on + // the GPU. So unless you seem to be actively working with the host side's + // data then we do this initialization on the device side since this avoids a + // host to device transfer that would likely immediately follow. + if (data().device_ready()) + { + cuda::set_tensor(*this, val); + return *this; + } + ) + auto d = host_write_only(); for (size_t i = 0; i < size(); ++i) d[i] = val; @@ -97,15 +110,16 @@ namespace dlib tensor& operator*= (float val) { -#ifdef DLIB_USE_CUDA - cuda::scale_tensor(*this, val); - return *this; -#else - for (auto& d : *this) - d *= val; + IF_DLIB_USE_CUDA( + cuda::scale_tensor(*this, val); + ) + + IF_DLIB_NOT_USE_CUDA( + for (auto& d : *this) + d *= val; + ) return *this; -#endif } tensor& operator/= (float val) diff --git a/dlib/cuda/tensor_tools.cpp b/dlib/cuda/tensor_tools.cpp index c6640b30e0..ce3d323c92 100644 --- a/dlib/cuda/tensor_tools.cpp +++ b/dlib/cuda/tensor_tools.cpp @@ -17,17 +17,6 @@ namespace dlib static std::atomic var(true); return var; } - - bool& use_cuda_impl ( - ) - { -#ifdef DLIB_USE_CUDA - thread_local bool var(cuda::is_available()); -#else - thread_local bool var(false); -#endif - return var; - } } bool dnn_prefer_fastest_algorithms ( @@ -47,21 +36,6 @@ namespace dlib { dnn_prefer_fastest_algo() = false; } - - bool use_cuda( - ) - { - return use_cuda_impl(); - } - - void set_use_cuda( - bool flag - ) - { -#ifdef DLIB_USE_CUDA - use_cuda_impl() = flag; -#endif - } } namespace dlib { namespace tt diff --git a/dlib/cuda/tensor_tools.h b/dlib/cuda/tensor_tools.h index f9ec1d11fe..e551555260 100644 --- a/dlib/cuda/tensor_tools.h +++ b/dlib/cuda/tensor_tools.h @@ -15,26 +15,11 @@ #include "../geometry/rectangle.h" #include "../test_for_odr_violations.h" -#ifdef DLIB_USE_CUDA -#define IF_DLIB_USE_CUDA(...) if (use_cuda()) { __VA_ARGS__ } -#else -#define IF_DLIB_USE_CUDA(...) -#endif - -#ifdef DLIB_USE_CUDA -#define IF_DLIB_NOT_USE_CUDA(...) if (!use_cuda()) { __VA_ARGS__ } -#else -#define IF_DLIB_NOT_USE_CUDA(...) __VA_ARGS__ -#endif - namespace dlib { bool dnn_prefer_fastest_algorithms(); void set_dnn_prefer_fastest_algorithms(); void set_dnn_prefer_smallest_algorithms(); - - bool use_cuda(); - void set_use_cuda(bool flag); } namespace dlib { namespace tt diff --git a/dlib/dnn/loss.h b/dlib/dnn/loss.h index 36b37a2956..6a8f257bec 100644 --- a/dlib/dnn/loss.h +++ b/dlib/dnn/loss.h @@ -2823,11 +2823,13 @@ namespace dlib } double loss; -#ifdef DLIB_USE_CUDA - cuda_compute(truth, output_tensor, grad, loss); -#else - cpu_compute(truth, output_tensor, grad, loss); -#endif + IF_DLIB_USE_CUDA( + cuda_compute(truth, output_tensor, grad, loss); + ) + + IF_DLIB_NOT_USE_CUDA( + cpu_compute(truth, output_tensor, grad, loss); + ) return loss; } @@ -2859,9 +2861,8 @@ namespace dlib #ifdef DLIB_USE_CUDA cuda::compute_loss_binary_log_per_pixel cuda_compute; -#else - cpu::compute_loss_binary_log_per_pixel cpu_compute; #endif + cpu::compute_loss_binary_log_per_pixel cpu_compute; }; template @@ -2968,11 +2969,13 @@ namespace dlib double loss; -#ifdef DLIB_USE_CUDA - cuda_compute(truth, output_tensor, grad, loss); -#else - cpu_compute(truth, output_tensor, grad, loss); -#endif + IF_DLIB_USE_CUDA( + cuda_compute(truth, output_tensor, grad, loss); + ) + + IF_DLIB_NOT_USE_CUDA( + cpu_compute(truth, output_tensor, grad, loss); + ) return loss; } @@ -3004,9 +3007,8 @@ namespace dlib #ifdef DLIB_USE_CUDA cuda::compute_loss_multiclass_log_per_pixel cuda_compute; -#else - cpu::compute_loss_multiclass_log_per_pixel cpu_compute; #endif + cpu::compute_loss_multiclass_log_per_pixel cpu_compute; }; template @@ -3068,11 +3070,13 @@ namespace dlib } double loss; -#ifdef DLIB_USE_CUDA - cuda_compute(truth, output_tensor, grad, loss); -#else - cpu_compute(truth, output_tensor, grad, loss); -#endif + IF_DLIB_USE_CUDA( + cuda_compute(truth, output_tensor, grad, loss); + ) + + IF_DLIB_NOT_USE_CUDA( + cpu_compute(truth, output_tensor, grad, loss); + ) return loss; } @@ -3104,9 +3108,8 @@ namespace dlib #ifdef DLIB_USE_CUDA cuda::compute_loss_multiclass_log_per_pixel_weighted cuda_compute; -#else - cpu::compute_loss_multiclass_log_per_pixel_weighted cpu_compute; #endif + cpu::compute_loss_multiclass_log_per_pixel_weighted cpu_compute; }; @@ -3319,11 +3322,13 @@ namespace dlib } } double loss; -#ifdef DLIB_USE_CUDA - cuda_compute(truth, output_tensor, grad, loss); -#else - cpu_compute(truth, output_tensor, grad, loss); -#endif + IF_DLIB_USE_CUDA( + cuda_compute(truth, output_tensor, grad, loss); + ) + + IF_DLIB_NOT_USE_CUDA( + cpu_compute(truth, output_tensor, grad, loss); + ) return loss; } @@ -3355,9 +3360,8 @@ namespace dlib #ifdef DLIB_USE_CUDA cuda::compute_loss_mean_squared_per_channel_and_pixel cuda_compute; -#else - cpu::compute_loss_mean_squared_per_channel_and_pixel cpu_compute; #endif + cpu::compute_loss_mean_squared_per_channel_and_pixel cpu_compute; }; template From 40dc2dd9d68b6cb2a553450221dfdd5beeab166d Mon Sep 17 00:00:00 2001 From: Davis King Date: Sat, 24 Jan 2026 19:05:12 -0500 Subject: [PATCH 25/39] update to use newer sphinx --- docs/docs/python/conf.py | 4 +++- docs/makedocs | 6 +++--- 2 files changed, 6 insertions(+), 4 deletions(-) diff --git a/docs/docs/python/conf.py b/docs/docs/python/conf.py index 7869d0273c..64143379ac 100644 --- a/docs/docs/python/conf.py +++ b/docs/docs/python/conf.py @@ -136,7 +136,9 @@ #html_use_smartypants = True # Custom sidebar templates, maps document names to template names. -#html_sidebars = {} +html_sidebars = { + "**": [] +} # Additional templates that should be rendered to pages, maps page names to # template names. diff --git a/docs/makedocs b/docs/makedocs index 872edf47cd..8cd0b17065 100755 --- a/docs/makedocs +++ b/docs/makedocs @@ -140,15 +140,15 @@ makedocs () then cd .. python setup.py build || report_failure - python setup.py build_sphinx -c docs/docs/python --build-dir docs/sphinx.$$ || report_failure + python -m sphinx -b html -c docs/docs/python docs/docs/python docs/sphinx.$$ || report_failure # sphinx will read in the _dlib_pybind11 module and use that to name everything. But that's # not what we want, so we rename that to dlib everywhere. You would think sphinx would be # able to deal with the dlib/__init__.py file and this wouldn't be necessary, but that # doesn't seem to be the case. find docs/sphinx.$$ -type f | xargs sed -i -e "s/_dlib_pybind11/dlib/g" cd docs - cp -r sphinx.$$/html docs/web/python - mv sphinx.$$/html docs/chm/docs/python + cp -r sphinx.$$ docs/web/python + mv sphinx.$$ docs/chm/docs/python rm -rf sphinx.$$ fi; From 5a490e0f50fc41082373f546f0f0e2cfbaf497e7 Mon Sep 17 00:00:00 2001 From: "Davis E. King" Date: Wed, 4 Feb 2026 19:48:30 -0500 Subject: [PATCH 26/39] update build rules to work with latest python build practices (#3134) * update build rules to work with latest python build practices * try this * add cmake --- .github/workflows/build_python.yml | 9 +++------ pyproject.toml | 2 +- 2 files changed, 4 insertions(+), 7 deletions(-) diff --git a/.github/workflows/build_python.yml b/.github/workflows/build_python.yml index 9860b9c7fc..d40b261c3e 100644 --- a/.github/workflows/build_python.yml +++ b/.github/workflows/build_python.yml @@ -44,8 +44,7 @@ jobs: - name: Build run: | pip3 install cmake==3.24.0 - python setup.py build - python setup.py install --user + pip3 install . - name: Test run: python -m pytest --ignore docs --ignore dlib @@ -58,8 +57,7 @@ jobs: run: pip install pytest numpy - name: Build run: | - python setup.py build - python setup.py install --user + pip install . - name: Test run: python -m pytest --ignore docs --ignore dlib @@ -73,8 +71,7 @@ jobs: # run: pip3 install pytest numpy # - name: Build # run: | -# python3 setup.py build -# python3 setup.py install --user +# pip install . # - name: Test # run: python3 -m pytest --ignore docs --ignore dlib diff --git a/pyproject.toml b/pyproject.toml index 2a301221f0..bf3839967b 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -1,4 +1,4 @@ [build-system] -requires = ["setuptools", "wheel", "packaging"] +requires = ["setuptools", "wheel", "packaging", "cmake"] build-backend = "setuptools.build_meta" From 94553fdc7964cedef340eced5b2f67896937ff00 Mon Sep 17 00:00:00 2001 From: Kane Scipioni Date: Sat, 14 Feb 2026 19:41:38 -0600 Subject: [PATCH 27/39] Update path to mkl and kiss fft headers (#3136) --- dlib/test/fft.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/dlib/test/fft.cpp b/dlib/test/fft.cpp index 43cbc6f89f..e2ff82a3bb 100644 --- a/dlib/test/fft.cpp +++ b/dlib/test/fft.cpp @@ -12,8 +12,8 @@ #include #ifdef DLIB_USE_MKL_FFT -#include -#include +#include +#include #endif #include "tester.h" #include "fftr_good_data.h" From 0e667c1c802a3cb67b78778f8874823cd8b9d566 Mon Sep 17 00:00:00 2001 From: Davis King Date: Mon, 9 Mar 2026 14:28:34 -0400 Subject: [PATCH 28/39] Improve numerical robustness of find_min_trust_region() --- dlib/optimization/optimization_trust_region.h | 10 ++- dlib/test/trust_region.cpp | 75 +++++++++++++++++++ 2 files changed, 83 insertions(+), 2 deletions(-) diff --git a/dlib/optimization/optimization_trust_region.h b/dlib/optimization/optimization_trust_region.h index 5f0ad897f9..9e9c62e242 100644 --- a/dlib/optimization/optimization_trust_region.h +++ b/dlib/optimization/optimization_trust_region.h @@ -453,8 +453,14 @@ namespace dlib // predicted_improvement shouldn't be negative but it might be if something went // wrong in the trust region solver. So put abs() here to guard against that. This // way the sign of rho is determined only by the sign of measured_improvement. - const type rho = measured_improvement/std::abs(predicted_improvement); - + // + // When the predicted improvement is below the floating point resolution of f_value, + // the measured improvement (f_value - new_f_value) is dominated by rounding noise + // and rho becomes meaningless. In this case we trust the quadratic model and treat + // the step as a perfect match (rho = 1). + const type abs_pred = std::abs(predicted_improvement); + const type f_eps = std::abs(f_value) * std::numeric_limits::epsilon(); + const type rho = (abs_pred < f_eps) ? type(1) : measured_improvement/abs_pred; if (!is_finite(rho)) break; diff --git a/dlib/test/trust_region.cpp b/dlib/test/trust_region.cpp index aa2775b9c9..b5f50c70f4 100644 --- a/dlib/test/trust_region.cpp +++ b/dlib/test/trust_region.cpp @@ -75,6 +75,79 @@ namespace DLIB_TEST_MSG(length(p-ans) < 1e-5, "length(p): " << length(p-ans)); } +// ---------------------------------------------------------------------------------------- + + // A model where the objective function differs from the quadratic model by a + // large constant offset. The Hessian and gradient reported to the trust region + // solver are exact (they describe the quadratic part), but the objective value + // returned by operator() includes a huge offset. This means the quadratic model + // predictions are perfect, yet the floating-point subtraction + // measured_improvement = f(x) - f(x+p) + // loses all significant digits once the true improvement drops below + // |offset| * machine_epsilon. + // + // To make the optimizer take many small steps (so it actually reaches the regime + // where rounding kills the measured improvement), we use a Rosenbrock-like + // function which is not spherically symmetric. The narrow curved valley forces + // the trust region to take many constrained steps that only slowly approach the + // minimum. + template + struct offset_rosen_model + { + typedef matrix column_vector; + typedef matrix general_matrix; + + const T offset; + explicit offset_rosen_model(const T off) : offset(off) {} + + T operator()(const column_vector& x) const + { + return static_cast(rosen(x)) + offset; + } + + void get_derivative_and_hessian( + const column_vector& x, + column_vector& d, + general_matrix& h + ) const + { + d = rosen_derivative(x); + h = rosen_hessian(x); + } + }; + + void test_rho_with_large_offset() + { + print_spinner(); + + // Rosenbrock + huge offset. The offset makes the floating-point + // subtraction f(x) - f(x+p) lose precision once the true improvement + // drops below |offset| * epsilon ≈ 0.22. Rosenbrock's narrow valley + // forces many small trust region steps, so the optimizer inevitably + // enters this low-improvement regime. Without the rho fix, the + // optimizer stalls far from the minimum because every step in that + // regime produces rho ≈ 0 (or garbage), shrinking the radius to zero. + const double offset = 1e15; + + matrix x; + x = -1.2, 1.0; + + const double result = find_min_trust_region( + objective_delta_stop_strategy(0, 500), + offset_rosen_model(offset), + x); + + matrix ans; + ans = 1, 1; + + dlog << LINFO << "offset rosen obj: " << result - offset; + dlog << LINFO << "offset rosen x: " << trans(x); + dlog << LINFO << "offset rosen error: " << length(x - ans); + + DLIB_TEST_MSG(length(x - ans) < 1e-4, + "optimizer failed to converge with large offset, error = " << length(x - ans)); + } + // ---------------------------------------------------------------------------------------- void test_trust_region_sub_problem() @@ -320,6 +393,8 @@ namespace test_trust_region_sub_problem(); + test_rho_with_large_offset(); + test_problems(); } } a; From adf759362c3cab954144a38a23c0f4e9fa987944 Mon Sep 17 00:00:00 2001 From: Alex Shek <82931936+ykshek@users.noreply.github.com> Date: Wed, 18 Mar 2026 06:56:17 +0800 Subject: [PATCH 29/39] fix(test/string): fix gcc 16 build issue (#3137) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * fix(test/string): avoid use of deleted function Remove portion of code to fix a build issue for gcc16(C++23), where the build fails due to test `error: use of deleted function ‘std::basic_ostream& std::operator<<(basic_ostream&, const wchar_t*) [with _Traits = char_traits]’ 148 | DLIB_TEST_MSG( (left_substr(ws,L".") == L"file"), L"");` * fix(test/tokenizer): remove u8"..." C++20 fails due to convert from ‘’ to ‘std::vector >’ * Revert "fix(test/string): avoid use of deleted function" This reverts commit ad56f130cf30e8940c8012d5c4b2ebce4b5af711. * fix(test/string): Attempt fix again without deleting the test --- dlib/test/string.cpp | 4 ++-- dlib/test/tokenizer.cpp | 8 ++++---- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/dlib/test/string.cpp b/dlib/test/string.cpp index 18f9035c5f..09ac3e304b 100644 --- a/dlib/test/string.cpp +++ b/dlib/test/string.cpp @@ -145,9 +145,9 @@ namespace wstring ws = L"file.txt"; DLIB_TEST( (left_substr(ws,wstring(L".")) == L"file")); - DLIB_TEST_MSG( (left_substr(ws,L".") == L"file"), L""); + DLIB_TEST_MSG( (left_substr(ws,L".") == L"file"), ""); DLIB_TEST( (right_substr(ws,wstring(L".")) == L"txt")); - DLIB_TEST_MSG( (right_substr(ws,L".") == L"txt"), L""); + DLIB_TEST_MSG( (right_substr(ws,L".") == L"txt"), ""); dlog << LTRACE << 8; diff --git a/dlib/test/tokenizer.cpp b/dlib/test/tokenizer.cpp index f4300666a2..24d24906ea 100644 --- a/dlib/test/tokenizer.cpp +++ b/dlib/test/tokenizer.cpp @@ -394,10 +394,10 @@ namespace deserialize(loaded_test, in_stream); std::vector test_strings = { - u8"This is a test of the tokenisation process...\nimplemented in the Dlib library!", // English - u8"Ceci est un test du processus de\ntokenisation implémenté dans\nla bibliothèque Dlib!", // French - u8"Dette er en test af tokeniseringsprocessen implementeret i Dlib-biblioteket!", // Danish - u8"这是对Dlib库中实现的标记化过程的测试！" // Chinese + "This is a test of the tokenisation process...\nimplemented in the Dlib library!", // English + "Ceci est un test du processus de\ntokenisation implémenté dans\nla bibliothèque Dlib!", // French + "Dette er en test af tokeniseringsprocessen implementeret i Dlib-biblioteket!", // Danish + "这是对Dlib库中实现的标记化过程的测试！" // Chinese }; for (const auto& text : test_strings) { From 1fdece330b9b3a02e9cf6c9b080eb62ed4c882db Mon Sep 17 00:00:00 2001 From: Davis King Date: Sun, 29 Mar 2026 13:05:49 -0400 Subject: [PATCH 30/39] tag 20.0.1 --- dlib/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dlib/CMakeLists.txt b/dlib/CMakeLists.txt index 04b5d92059..0c7f43bf18 100644 --- a/dlib/CMakeLists.txt +++ b/dlib/CMakeLists.txt @@ -18,7 +18,7 @@ project(dlib LANGUAGES C CXX) set(CPACK_PACKAGE_NAME "dlib") set(CPACK_PACKAGE_VERSION_MAJOR "20") set(CPACK_PACKAGE_VERSION_MINOR "0") -set(CPACK_PACKAGE_VERSION_PATCH "99") +set(CPACK_PACKAGE_VERSION_PATCH "1") set(VERSION ${CPACK_PACKAGE_VERSION_MAJOR}.${CPACK_PACKAGE_VERSION_MINOR}.${CPACK_PACKAGE_VERSION_PATCH}) # Only print these messages once, even if dlib is added multiple times via add_subdirectory() if (NOT TARGET dlib) From fcf5fbea9ee6164534b721d58e9a99590c5c0ce8 Mon Sep 17 00:00:00 2001 From: Davis King Date: Sun, 29 Mar 2026 13:06:42 -0400 Subject: [PATCH 31/39] set back to .99 --- dlib/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dlib/CMakeLists.txt b/dlib/CMakeLists.txt index 0c7f43bf18..04b5d92059 100644 --- a/dlib/CMakeLists.txt +++ b/dlib/CMakeLists.txt @@ -18,7 +18,7 @@ project(dlib LANGUAGES C CXX) set(CPACK_PACKAGE_NAME "dlib") set(CPACK_PACKAGE_VERSION_MAJOR "20") set(CPACK_PACKAGE_VERSION_MINOR "0") -set(CPACK_PACKAGE_VERSION_PATCH "1") +set(CPACK_PACKAGE_VERSION_PATCH "99") set(VERSION ${CPACK_PACKAGE_VERSION_MAJOR}.${CPACK_PACKAGE_VERSION_MINOR}.${CPACK_PACKAGE_VERSION_PATCH}) # Only print these messages once, even if dlib is added multiple times via add_subdirectory() if (NOT TARGET dlib) From eea87b98de8074042f57f7362da82a2dcc3e85b1 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Yasin=20G=C3=96Z=C3=9CB=C3=9CY=C3=9CK?= Date: Sat, 25 Apr 2026 19:12:52 +0300 Subject: [PATCH 32/39] Add undo and redo functionality to imglab (#606) (#3143) Implemented Ctrl+Z and Ctrl+Y/Ctrl+Shift+Z shortcuts in imglab to undo and redo bounding box modifications. Added equality operators to box struct and limited maximum state history to 10 to avoid large memory footprints. --- dlib/data_io/image_dataset_metadata.h | 21 +++++++ tools/imglab/src/metadata_editor.cpp | 79 +++++++++++++++++++++++++++ tools/imglab/src/metadata_editor.h | 7 +++ 3 files changed, 107 insertions(+) diff --git a/dlib/data_io/image_dataset_metadata.h b/dlib/data_io/image_dataset_metadata.h index 7f3ede426a..acfbff90ac 100644 --- a/dlib/data_io/image_dataset_metadata.h +++ b/dlib/data_io/image_dataset_metadata.h @@ -93,6 +93,27 @@ namespace dlib ensures - returns true if label metadata is present and false otherwise. !*/ + + bool operator==(const box& rhs) const + { + return rect == rhs.rect && + parts == rhs.parts && + label == rhs.label && + difficult == rhs.difficult && + truncated == rhs.truncated && + occluded == rhs.occluded && + ignore == rhs.ignore && + pose == rhs.pose && + detection_score == rhs.detection_score && + angle == rhs.angle && + gender == rhs.gender && + age == rhs.age; + } + + bool operator!=(const box& rhs) const + { + return !(*this == rhs); + } }; // ------------------------------------------------------------------------------------ diff --git a/tools/imglab/src/metadata_editor.cpp b/tools/imglab/src/metadata_editor.cpp index f3e0e49db2..459f1fb88c 100644 --- a/tools/imglab/src/metadata_editor.cpp +++ b/tools/imglab/src/metadata_editor.cpp @@ -21,6 +21,11 @@ extern const char* VERSION; // ---------------------------------------------------------------------------------------- +std::vector get_overlays ( + const dlib::image_dataset_metadata::image& data, + color_mapper& string_to_color +); + metadata_editor:: metadata_editor( const std::string& filename_, @@ -390,6 +395,25 @@ on_keydown ( select_image(image_pos); } + if ((key == 'z' || key == 'Z') && (state&base_window::KBD_MOD_CONTROL) && !overlay_label.has_input_focus()) + { + if (state&base_window::KBD_MOD_SHIFT) + { + perform_redo(); + } + else + { + perform_undo(); + } + return; + } + + if ((key == 'y' || key == 'Y') && (state&base_window::KBD_MOD_CONTROL) && !overlay_label.has_input_focus()) + { + perform_redo(); + return; + } + // Make 'w' and 's' act like KEY_UP and KEY_DOWN if ((key == 'w' || key == 'W') && !overlay_label.has_input_focus()) { @@ -532,6 +556,12 @@ load_image( if (idx >= metadata.images.size()) return; + if (image_pos != idx) + { + undo_history.clear(); + redo_history.clear(); + } + image_pos = idx; array2d img; @@ -556,6 +586,40 @@ load_image( // ---------------------------------------------------------------------------------------- +void metadata_editor:: +perform_undo() +{ + if (!undo_history.empty()) + { + redo_history.push_back(metadata.images[image_pos].boxes); + if (redo_history.size() > MAX_UNDO_HISTORY) redo_history.erase(redo_history.begin()); + + metadata.images[image_pos].boxes = undo_history.back(); + undo_history.pop_back(); + display.clear_overlay(); + display.add_overlay(get_overlays(metadata.images[image_pos], string_to_color)); + } +} + +// ---------------------------------------------------------------------------------------- + +void metadata_editor:: +perform_redo() +{ + if (!redo_history.empty()) + { + undo_history.push_back(metadata.images[image_pos].boxes); + if (undo_history.size() > MAX_UNDO_HISTORY) undo_history.erase(undo_history.begin()); + + metadata.images[image_pos].boxes = redo_history.back(); + redo_history.pop_back(); + display.clear_overlay(); + display.add_overlay(get_overlays(metadata.images[image_pos], string_to_color)); + } +} + +// ---------------------------------------------------------------------------------------- + void metadata_editor:: load_image_and_set_size( unsigned long idx @@ -564,6 +628,12 @@ load_image_and_set_size( if (idx >= metadata.images.size()) return; + if (image_pos != idx) + { + undo_history.clear(); + redo_history.clear(); + } + image_pos = idx; array2d img; @@ -616,6 +686,7 @@ on_overlay_rects_changed( const std::vector& rects = display.get_overlay_rects(); std::vector& boxes = metadata.images[image_pos].boxes; + std::vector old_boxes = boxes; boxes.clear(); for (unsigned long i = 0; i < rects.size(); ++i) @@ -627,6 +698,14 @@ on_overlay_rects_changed( temp.ignore = rects[i].crossed_out; boxes.push_back(temp); } + + if (old_boxes != boxes) + { + undo_history.push_back(old_boxes); + if (undo_history.size() > MAX_UNDO_HISTORY) + undo_history.erase(undo_history.begin()); + redo_history.clear(); + } } } diff --git a/tools/imglab/src/metadata_editor.h b/tools/imglab/src/metadata_editor.h index eba3ed3a44..6320e36256 100644 --- a/tools/imglab/src/metadata_editor.h +++ b/tools/imglab/src/metadata_editor.h @@ -7,6 +7,7 @@ #include #include #include +#include // ---------------------------------------------------------------------------------------- @@ -69,6 +70,8 @@ class metadata_editor : public dlib::drawable_window ); private: + void perform_undo(); + void perform_redo(); void file_save(); void file_save_as(); @@ -96,6 +99,10 @@ class metadata_editor : public dlib::drawable_window std::string filename; dlib::image_dataset_metadata::dataset metadata; + std::vector> undo_history; + std::vector> redo_history; + static constexpr size_t MAX_UNDO_HISTORY = 10; + dlib::menu_bar mbar; dlib::list_box lb_images; unsigned long image_pos; From dcc211e402b3780614581b66e118456f1e74cbf3 Mon Sep 17 00:00:00 2001 From: Davis King Date: Sat, 25 Apr 2026 12:13:07 -0400 Subject: [PATCH 33/39] Do some cleanup --- dlib/data_io/image_dataset_metadata.cpp | 32 ++++++++++- dlib/data_io/image_dataset_metadata.h | 37 ++++++------- dlib/test/data_io.cpp | 72 ++++++++++++++++++++++++- tools/imglab/src/metadata_editor.cpp | 48 ++++++++++++----- tools/imglab/src/metadata_editor.h | 11 ++-- 5 files changed, 159 insertions(+), 41 deletions(-) diff --git a/dlib/data_io/image_dataset_metadata.cpp b/dlib/data_io/image_dataset_metadata.cpp index e1ca410898..76dffb8a7c 100644 --- a/dlib/data_io/image_dataset_metadata.cpp +++ b/dlib/data_io/image_dataset_metadata.cpp @@ -19,6 +19,37 @@ namespace dlib namespace image_dataset_metadata { + // ------------------------------------------------------------------------------------ + + bool box:: + operator== ( + const box& rhs + ) const + { + return rect == rhs.rect && + parts == rhs.parts && + label == rhs.label && + difficult == rhs.difficult && + truncated == rhs.truncated && + occluded == rhs.occluded && + ignore == rhs.ignore && + pose == rhs.pose && + detection_score == rhs.detection_score && + angle == rhs.angle && + gender == rhs.gender && + age == rhs.age; + } + + // ------------------------------------------------------------------------------------ + + bool box:: + operator!= ( + const box& rhs + ) const + { + return !(*this == rhs); + } + // ------------------------------------------------------------------------------------ const std::string get_decoded_string(); @@ -417,4 +448,3 @@ namespace dlib #endif // DLIB_IMAGE_DAtASET_METADATA_CPPh_ - diff --git a/dlib/data_io/image_dataset_metadata.h b/dlib/data_io/image_dataset_metadata.h index acfbff90ac..1f868da6c4 100644 --- a/dlib/data_io/image_dataset_metadata.h +++ b/dlib/data_io/image_dataset_metadata.h @@ -94,26 +94,22 @@ namespace dlib - returns true if label metadata is present and false otherwise. !*/ - bool operator==(const box& rhs) const - { - return rect == rhs.rect && - parts == rhs.parts && - label == rhs.label && - difficult == rhs.difficult && - truncated == rhs.truncated && - occluded == rhs.occluded && - ignore == rhs.ignore && - pose == rhs.pose && - detection_score == rhs.detection_score && - angle == rhs.angle && - gender == rhs.gender && - age == rhs.age; - } - - bool operator!=(const box& rhs) const - { - return !(*this == rhs); - } + bool operator== ( + const box& rhs + ) const; + /*! + ensures + - returns true if and only if all member variables in *this compare + equal to their corresponding member variables in rhs. + !*/ + + bool operator!= ( + const box& rhs + ) const; + /*! + ensures + - returns !(*this == rhs) + !*/ }; // ------------------------------------------------------------------------------------ @@ -194,4 +190,3 @@ namespace dlib #endif #endif // DLIB_IMAGE_DAtASET_METADATA_Hh_ - diff --git a/dlib/test/data_io.cpp b/dlib/test/data_io.cpp index 8ced880084..550e249f00 100644 --- a/dlib/test/data_io.cpp +++ b/dlib/test/data_io.cpp @@ -204,6 +204,75 @@ namespace } } + void test_image_dataset_metadata_box_equality() + { + using namespace image_dataset_metadata; + + box a(rectangle(1,2,3,4)); + a.parts["nose"] = point(2,3); + a.label = "person"; + a.difficult = true; + a.truncated = true; + a.occluded = true; + a.ignore = true; + a.pose = 1.5; + a.detection_score = 0.75; + a.angle = 0.25; + a.gender = FEMALE; + a.age = 34.5; + + box b = a; + DLIB_TEST(a == b); + DLIB_TEST(!(a != b)); + + b.rect = rectangle(2,3,4,5); + DLIB_TEST(a != b); + + b = a; + b.parts["left_eye"] = point(1,2); + DLIB_TEST(a != b); + + b = a; + b.label = "car"; + DLIB_TEST(a != b); + + b = a; + b.difficult = false; + DLIB_TEST(a != b); + + b = a; + b.truncated = false; + DLIB_TEST(a != b); + + b = a; + b.occluded = false; + DLIB_TEST(a != b); + + b = a; + b.ignore = false; + DLIB_TEST(a != b); + + b = a; + b.pose = 2.5; + DLIB_TEST(a != b); + + b = a; + b.detection_score = 0.5; + DLIB_TEST(a != b); + + b = a; + b.angle = 0.5; + DLIB_TEST(a != b); + + b = a; + b.gender = MALE; + DLIB_TEST(a != b); + + b = a; + b.age = 42.0; + DLIB_TEST(a != b); + } + void perform_test ( ) @@ -211,6 +280,7 @@ namespace print_spinner(); create_iris_datafile(); + test_image_dataset_metadata_box_equality(); test_sparse_to_dense(); run_test >(); @@ -223,5 +293,3 @@ namespace test_data_io a; } - - diff --git a/tools/imglab/src/metadata_editor.cpp b/tools/imglab/src/metadata_editor.cpp index 459f1fb88c..850b780cdc 100644 --- a/tools/imglab/src/metadata_editor.cpp +++ b/tools/imglab/src/metadata_editor.cpp @@ -20,6 +20,7 @@ using namespace dlib; extern const char* VERSION; // ---------------------------------------------------------------------------------------- +constexpr size_t MAX_UNDO_HISTORY = 10; std::vector get_overlays ( const dlib::image_dataset_metadata::image& data, @@ -176,6 +177,8 @@ file_save_as() void metadata_editor:: remove_selected_images() { + clear_undo_history(); + dlib::queue::kernel_1a list; lb_images.get_selected(list); list.reset(); @@ -366,8 +369,7 @@ on_keydown ( if (keyboard_jump_pos >= metadata.images.size()) keyboard_jump_pos = metadata.images.size()-1; - image_pos = keyboard_jump_pos; - select_image(image_pos); + select_image(keyboard_jump_pos); } else { @@ -558,8 +560,7 @@ load_image( if (image_pos != idx) { - undo_history.clear(); - redo_history.clear(); + clear_undo_history(); } image_pos = idx; @@ -586,13 +587,36 @@ load_image( // ---------------------------------------------------------------------------------------- +void metadata_editor:: +clear_undo_history() +{ + undo_history.clear(); + redo_history.clear(); +} + +// ---------------------------------------------------------------------------------------- + +void metadata_editor:: +save_undo_state ( + const box_history_entry& boxes +) +{ + undo_history.push_back(boxes); + if (undo_history.size() > MAX_UNDO_HISTORY) + undo_history.erase(undo_history.begin()); + redo_history.clear(); +} + +// ---------------------------------------------------------------------------------------- + void metadata_editor:: perform_undo() { if (!undo_history.empty()) { redo_history.push_back(metadata.images[image_pos].boxes); - if (redo_history.size() > MAX_UNDO_HISTORY) redo_history.erase(redo_history.begin()); + if (redo_history.size() > MAX_UNDO_HISTORY) + redo_history.erase(redo_history.begin()); metadata.images[image_pos].boxes = undo_history.back(); undo_history.pop_back(); @@ -609,7 +633,8 @@ perform_redo() if (!redo_history.empty()) { undo_history.push_back(metadata.images[image_pos].boxes); - if (undo_history.size() > MAX_UNDO_HISTORY) undo_history.erase(undo_history.begin()); + if (undo_history.size() > MAX_UNDO_HISTORY) + undo_history.erase(undo_history.begin()); metadata.images[image_pos].boxes = redo_history.back(); redo_history.pop_back(); @@ -630,8 +655,7 @@ load_image_and_set_size( if (image_pos != idx) { - undo_history.clear(); - redo_history.clear(); + clear_undo_history(); } image_pos = idx; @@ -701,10 +725,7 @@ on_overlay_rects_changed( if (old_boxes != boxes) { - undo_history.push_back(old_boxes); - if (undo_history.size() > MAX_UNDO_HISTORY) - undo_history.erase(undo_history.begin()); - redo_history.clear(); + save_undo_state(old_boxes); } } } @@ -777,6 +798,8 @@ display_about( "and drag allows you to navigate around the image. Holding ctrl and " "left clicking a rectangle will give it the label from the Next Label field. " "Holding shift + right click and then dragging allows you to move things around. " + "Pressing ctrl+z will undo changes to the current image's boxes and pressing " + "ctrl+y or ctrl+shift+z will redo them. " "Holding ctrl and pressing the up or down keyboard keys will propagate " "rectangle labels from one image to the next and also skip empty images. " "Similarly, holding ctrl+shift will propagate entire boxes via a visual tracking " @@ -791,4 +814,3 @@ display_about( } // ---------------------------------------------------------------------------------------- - diff --git a/tools/imglab/src/metadata_editor.h b/tools/imglab/src/metadata_editor.h index 6320e36256..c1e833d7b7 100644 --- a/tools/imglab/src/metadata_editor.h +++ b/tools/imglab/src/metadata_editor.h @@ -72,6 +72,10 @@ class metadata_editor : public dlib::drawable_window private: void perform_undo(); void perform_redo(); + void clear_undo_history(); + void save_undo_state ( + const std::vector& boxes + ); void file_save(); void file_save_as(); @@ -99,9 +103,9 @@ class metadata_editor : public dlib::drawable_window std::string filename; dlib::image_dataset_metadata::dataset metadata; - std::vector> undo_history; - std::vector> redo_history; - static constexpr size_t MAX_UNDO_HISTORY = 10; + typedef std::vector box_history_entry; + std::vector undo_history; + std::vector redo_history; dlib::menu_bar mbar; dlib::list_box lb_images; @@ -123,4 +127,3 @@ class metadata_editor : public dlib::drawable_window #endif // DLIB_METADATA_EdITOR_H__ - From d1b6bcce2994eb2ae71687dad803cb3af8ec8b8f Mon Sep 17 00:00:00 2001 From: Davis King Date: Sat, 25 Apr 2026 13:46:50 -0400 Subject: [PATCH 34/39] switch to a single immutable global env var that controls cuda use --- dlib/cuda/cuda_dlib.cu | 56 ++++++++++++++++++++++--------- dlib/cuda/cuda_dlib.h | 9 ----- dlib/cuda/curand_dlibapi.cpp | 3 +- dlib/cuda/gpu_data.cpp | 22 +++++-------- dlib/cuda/tensor_tools.cpp | 64 ++++++++++++++++++++---------------- dlib/dnn/core_abstract.h | 36 ++++++++------------ 6 files changed, 100 insertions(+), 90 deletions(-) diff --git a/dlib/cuda/cuda_dlib.cu b/dlib/cuda/cuda_dlib.cu index d58681306f..9ff316c799 100644 --- a/dlib/cuda/cuda_dlib.cu +++ b/dlib/cuda/cuda_dlib.cu @@ -5,16 +5,37 @@ #include "cuda_dlib.h" #include "cudnn_dlibapi.h" #include +#include +#include namespace dlib { namespace { - bool& use_cuda_impl ( + bool cuda_device_available ( ) { - thread_local bool var(cuda::is_available()); + int num_devices; + return cudaGetDeviceCount(&num_devices) == cudaSuccess && num_devices > 0; + } + + bool cuda_disabled_by_environment ( + ) + { + const char* var = std::getenv("DLIB_DISABLE_CUDA_USE"); + return var != nullptr && + std::strcmp(var, "") != 0 && + std::strcmp(var, "0") != 0 && + std::strcmp(var, "false") != 0 && + std::strcmp(var, "False") != 0 && + std::strcmp(var, "FALSE") != 0; + } + + bool use_cuda_impl ( + ) + { + static const bool var = !cuda_disabled_by_environment() && cuda_device_available(); return var; } @@ -29,7 +50,7 @@ namespace dlib int dev ) { - if (is_available()) + if (use_cuda()) CHECK_CUDA(cudaSetDevice(dev)); } @@ -37,7 +58,7 @@ namespace dlib ) { int dev = -1; - if (is_available()) + if (use_cuda()) CHECK_CUDA(cudaGetDevice(&dev)); return dev; } @@ -46,6 +67,9 @@ namespace dlib int device ) { + if (!use_cuda()) + return "CUDA_DISABLED"; + cudaDeviceProp props; CHECK_CUDA(cudaGetDeviceProperties(&props, device)); return props.name; @@ -54,14 +78,14 @@ namespace dlib void set_current_device_blocking_sync( ) { - CHECK_CUDA(cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync)); + if (use_cuda()) + CHECK_CUDA(cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync)); } bool is_available( ) { - int num_devices; - return cudaGetDeviceCount(&num_devices) == cudaSuccess && num_devices > 0; + return use_cuda(); } bool use_cuda( @@ -70,17 +94,12 @@ namespace dlib return use_cuda_impl(); } - void set_use_cuda( - bool flag - ) - { - if (is_available()) - use_cuda_impl() = flag; - } - int get_num_devices ( ) { + if (!use_cuda()) + return 0; + int num_devices; CHECK_CUDA(cudaGetDeviceCount(&num_devices)); return num_devices; @@ -88,6 +107,9 @@ namespace dlib bool can_access_peer (int device_id, int peer_device_id) { + if (!use_cuda()) + return false; + int can_access; CHECK_CUDA(cudaDeviceCanAccessPeer(&can_access, device_id, peer_device_id)); return can_access != 0; @@ -99,6 +121,9 @@ namespace dlib void device_synchronize (int dev) { + if (!use_cuda()) + return; + raii_set_device set_dev(dev); CHECK_CUDA(cudaDeviceSynchronize()); } @@ -3254,4 +3279,3 @@ namespace dlib } } - diff --git a/dlib/cuda/cuda_dlib.h b/dlib/cuda/cuda_dlib.h index 8431413b9e..24f2a82a55 100644 --- a/dlib/cuda/cuda_dlib.h +++ b/dlib/cuda/cuda_dlib.h @@ -31,10 +31,6 @@ namespace dlib bool use_cuda( ); - void set_use_cuda( - bool flag - ); - std::string get_device_name ( int device ); @@ -958,10 +954,6 @@ namespace dlib inline bool use_cuda( ) { return false; } - inline void set_use_cuda( - bool flag - ) {} - inline std::string get_device_name ( int device ) @@ -999,4 +991,3 @@ namespace dlib #endif // DLIB_DNN_CuDA_H_ - diff --git a/dlib/cuda/curand_dlibapi.cpp b/dlib/cuda/curand_dlibapi.cpp index 6f8b00006b..02afed46d9 100644 --- a/dlib/cuda/curand_dlibapi.cpp +++ b/dlib/cuda/curand_dlibapi.cpp @@ -48,7 +48,7 @@ namespace dlib unsigned long long seed ) : handle(nullptr) { - if (is_available()) + if (use_cuda()) { curandGenerator_t gen; CHECK_CURAND(curandCreateGenerator(&gen, CURAND_RNG_PSEUDO_DEFAULT)); @@ -114,4 +114,3 @@ namespace dlib #endif // DLIB_USE_CUDA #endif // DLIB_DNN_CuRAND_CPP_ - diff --git a/dlib/cuda/gpu_data.cpp b/dlib/cuda/gpu_data.cpp index 413c3996ee..9c252bb7ea 100644 --- a/dlib/cuda/gpu_data.cpp +++ b/dlib/cuda/gpu_data.cpp @@ -54,20 +54,20 @@ namespace dlib } else { + if (!cuda::use_cuda()) + { + std::memcpy(dest.host()+dest_offset, src.host()+src_offset, num*sizeof(float)); + return; + } + // if we write to the entire thing then we can use device_write_only() if (dest_offset == 0 && num == dest.size()) { // copy the memory efficiently based on which copy is current in each object. - if (dest.device_id() >= 0 && src.device_ready()) + if (src.device_ready()) CHECK_CUDA(cudaMemcpy(dest.device_write_only(), src.device()+src_offset, num*sizeof(float), cudaMemcpyDeviceToDevice)); - else if (dest.device_id() < 0 && src.device_ready()) - CHECK_CUDA(cudaMemcpy(dest.host_write_only(), src.device()+src_offset, num*sizeof(float), cudaMemcpyDeviceToHost)); - else if (dest.device_id() >= 0 && !src.device_ready()) - CHECK_CUDA(cudaMemcpy(dest.device_write_only(), src.host()+src_offset, num*sizeof(float), cudaMemcpyHostToDevice)); - else if (dest.device_id() >= 0 || src.device_id() >= 0) - CHECK_CUDA(cudaMemcpy(dest.host_write_only(), src.host()+src_offset, num*sizeof(float), cudaMemcpyHostToHost)); else - std::memcpy(dest.host_write_only(), src.host()+src_offset, num*sizeof(float)); + CHECK_CUDA(cudaMemcpy(dest.device_write_only(), src.host()+src_offset, num*sizeof(float), cudaMemcpyHostToDevice)); } else { @@ -78,11 +78,8 @@ namespace dlib CHECK_CUDA(cudaMemcpy(dest.host()+dest_offset, src.device()+src_offset, num*sizeof(float), cudaMemcpyDeviceToHost)); else if (dest.device_ready() && !src.device_ready()) CHECK_CUDA(cudaMemcpy(dest.device()+dest_offset, src.host()+src_offset, num*sizeof(float), cudaMemcpyHostToDevice)); - else if (dest.device_id() >= 0 || src.device_id() >= 0) - CHECK_CUDA(cudaMemcpy(dest.host()+dest_offset, src.host()+src_offset, num*sizeof(float), cudaMemcpyHostToHost)); else - std::memcpy(dest.host()+dest_offset, src.host()+src_offset, num*sizeof(float)); - + CHECK_CUDA(cudaMemcpy(dest.host()+dest_offset, src.host()+src_offset, num*sizeof(float), cudaMemcpyHostToHost)); } } } @@ -267,4 +264,3 @@ namespace dlib #endif // DLIB_USE_CUDA #endif // DLIB_GPU_DaTA_CPP_ - diff --git a/dlib/cuda/tensor_tools.cpp b/dlib/cuda/tensor_tools.cpp index 69ccf9ebc4..be01d38533 100644 --- a/dlib/cuda/tensor_tools.cpp +++ b/dlib/cuda/tensor_tools.cpp @@ -1585,13 +1585,15 @@ namespace dlib { namespace tt long feature_dim ) { -#ifdef DLIB_USE_CUDA - cuda::compute_act_halt_probabilities(halt_probs, logits, input_data, halt_params, - batch_size, seq_len, feature_dim); -#else - cpu::compute_act_halt_probabilities(halt_probs, logits, input_data, halt_params, - batch_size, seq_len, feature_dim); -#endif + IF_DLIB_USE_CUDA( + cuda::compute_act_halt_probabilities(halt_probs, logits, input_data, halt_params, + batch_size, seq_len, feature_dim); + ) + + IF_DLIB_NOT_USE_CUDA( + cpu::compute_act_halt_probabilities(halt_probs, logits, input_data, halt_params, + batch_size, seq_len, feature_dim); + ) } void update_act_state( @@ -1610,13 +1612,15 @@ namespace dlib { namespace tt long current_step ) { -#ifdef DLIB_USE_CUDA - cuda::update_act_state(output, input_data, halt_probs, cumulative_halting, remainders, - n_steps, effective_weights, batch_size, seq_len, d_model, num_channels, halt_threshold, current_step); -#else - cpu::update_act_state(output, input_data, halt_probs, cumulative_halting, remainders, - n_steps, effective_weights, batch_size, seq_len, d_model, num_channels, halt_threshold, current_step); -#endif + IF_DLIB_USE_CUDA( + cuda::update_act_state(output, input_data, halt_probs, cumulative_halting, remainders, + n_steps, effective_weights, batch_size, seq_len, d_model, num_channels, halt_threshold, current_step); + ) + + IF_DLIB_NOT_USE_CUDA( + cpu::update_act_state(output, input_data, halt_probs, cumulative_halting, remainders, + n_steps, effective_weights, batch_size, seq_len, d_model, num_channels, halt_threshold, current_step); + ) } void finalize_act_output( @@ -1630,13 +1634,15 @@ namespace dlib { namespace tt long num_channels ) { -#ifdef DLIB_USE_CUDA - cuda::finalize_act_output(output, input_data, remainders, effective_weights, - batch_size, seq_len, d_model, num_channels); -#else - cpu::finalize_act_output(output, input_data, remainders, effective_weights, - batch_size, seq_len, d_model, num_channels); -#endif + IF_DLIB_USE_CUDA( + cuda::finalize_act_output(output, input_data, remainders, effective_weights, + batch_size, seq_len, d_model, num_channels); + ) + + IF_DLIB_NOT_USE_CUDA( + cpu::finalize_act_output(output, input_data, remainders, effective_weights, + batch_size, seq_len, d_model, num_channels); + ) } void apply_act_depth_scaling( @@ -1650,13 +1656,15 @@ namespace dlib { namespace tt float scale_factor ) { -#ifdef DLIB_USE_CUDA - cuda::apply_act_depth_scaling(gradients, n_steps, batch_size, seq_len, - d_model, num_channels, max_steps, scale_factor); -#else - cpu::apply_act_depth_scaling(gradients, n_steps, batch_size, seq_len, - d_model, num_channels, max_steps, scale_factor); -#endif + IF_DLIB_USE_CUDA( + cuda::apply_act_depth_scaling(gradients, n_steps, batch_size, seq_len, + d_model, num_channels, max_steps, scale_factor); + ) + + IF_DLIB_NOT_USE_CUDA( + cpu::apply_act_depth_scaling(gradients, n_steps, batch_size, seq_len, + d_model, num_channels, max_steps, scale_factor); + ) } // ---------------------------------------------------------------------------------------- diff --git a/dlib/dnn/core_abstract.h b/dlib/dnn/core_abstract.h index 9d61c716b3..8c61af5bf9 100644 --- a/dlib/dnn/core_abstract.h +++ b/dlib/dnn/core_abstract.h @@ -196,27 +196,20 @@ namespace dlib - #dnn_prefer_fastest_algorithms() == false !*/ - bool use_cuda( - ); - /*! - ensures - - If dlib should use the CUDA implementation of a deep neural network - then this function returns true and false otherwise. - - On program startup this function will return true if DLIB_USE_CUDA is defined and - there is an available GPU device to use. - - This function always returns false if DLIB_USE_CUDA is not defined. - - This function sets a thread local variable. That is, each thread has its own value - for use_cuda(). This means that one thread may use cuda while another thread might - not, depending on the setting of use_cuda(). - !*/ - - void set_use_cuda( - bool flag - ); - /*! - ensures - - #use_cuda() == flag for the calling thread. - !*/ + namespace cuda + { + bool use_cuda( + ); + /*! + ensures + - If dlib should use CUDA then this function returns true and false otherwise. + - This function returns true if DLIB_USE_CUDA is defined, + the DLIB_DISABLE_CUDA_USE environment variable is not set to a true value, + and there is an available GPU device to use. + - This function always returns false if DLIB_USE_CUDA is not defined. + - The value returned by this function is fixed after its first call. + !*/ + } // ---------------------------------------------------------------------------------------- @@ -2103,4 +2096,3 @@ namespace dlib } #endif // DLIB_DNn_CORE_ABSTRACT_H_ - From 5447b108bd3204eaac91d6b9b09879b5a0d456c1 Mon Sep 17 00:00:00 2001 From: Davis King Date: Sun, 3 May 2026 18:44:22 -0400 Subject: [PATCH 35/39] more cleanup --- dlib/cuda/cuda_dlib.cu | 7 +------ dlib/cuda/cuda_dlib.h | 7 +------ dlib/cuda/curand_dlibapi.cpp | 1 + dlib/cuda/gpu_data.cpp | 3 ++- dlib/cuda/gpu_data.h | 25 ++++++++++++++++--------- dlib/cuda/gpu_data_abstract.h | 14 +++++++++++--- dlib/cuda/tensor.h | 7 +++++-- dlib/cuda/tensor_abstract.h | 10 +++++++--- dlib/cuda/tensor_tools.cpp | 10 +++++----- dlib/dnn/core_abstract.h | 5 ++++- dlib/test/dnn.cpp | 2 +- 11 files changed, 54 insertions(+), 37 deletions(-) diff --git a/dlib/cuda/cuda_dlib.cu b/dlib/cuda/cuda_dlib.cu index 9ff316c799..477f917452 100644 --- a/dlib/cuda/cuda_dlib.cu +++ b/dlib/cuda/cuda_dlib.cu @@ -82,12 +82,6 @@ namespace dlib CHECK_CUDA(cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync)); } - bool is_available( - ) - { - return use_cuda(); - } - bool use_cuda( ) { @@ -3279,3 +3273,4 @@ namespace dlib } } + diff --git a/dlib/cuda/cuda_dlib.h b/dlib/cuda/cuda_dlib.h index 24f2a82a55..5463fe7e77 100644 --- a/dlib/cuda/cuda_dlib.h +++ b/dlib/cuda/cuda_dlib.h @@ -25,9 +25,6 @@ namespace dlib int get_num_devices ( ); - bool is_available ( - ); - bool use_cuda( ); @@ -948,9 +945,6 @@ namespace dlib inline int get_num_devices ( ) { return 1; } - inline bool is_available ( - ) { return false; } - inline bool use_cuda( ) { return false; } @@ -991,3 +985,4 @@ namespace dlib #endif // DLIB_DNN_CuDA_H_ + diff --git a/dlib/cuda/curand_dlibapi.cpp b/dlib/cuda/curand_dlibapi.cpp index 02afed46d9..1a77f3b8aa 100644 --- a/dlib/cuda/curand_dlibapi.cpp +++ b/dlib/cuda/curand_dlibapi.cpp @@ -114,3 +114,4 @@ namespace dlib #endif // DLIB_USE_CUDA #endif // DLIB_DNN_CuRAND_CPP_ + diff --git a/dlib/cuda/gpu_data.cpp b/dlib/cuda/gpu_data.cpp index 9c252bb7ea..17b85288f8 100644 --- a/dlib/cuda/gpu_data.cpp +++ b/dlib/cuda/gpu_data.cpp @@ -187,6 +187,7 @@ namespace dlib host_current = true; device_current = true; device_in_use = false; + the_device_id = 0; data_host.reset(); data_device.reset(); } @@ -208,7 +209,7 @@ namespace dlib if (!cuda::use_cuda()) { data_host.reset(new float[new_size], std::default_delete()); - the_device_id = -1; + the_device_id = 0; return; } diff --git a/dlib/cuda/gpu_data.h b/dlib/cuda/gpu_data.h index 4d37315a40..5b42bc0640 100644 --- a/dlib/cuda/gpu_data.h +++ b/dlib/cuda/gpu_data.h @@ -101,16 +101,14 @@ namespace dlib float* host() { copy_to_host(); - if (device_id() >= 0) - device_current = false; + device_current = false; return data_host.get(); } float* host_write_only() { host_current = true; - if (device_id() >= 0) - device_current = false; + device_current = false; return data_host.get(); } @@ -118,8 +116,9 @@ namespace dlib { #ifndef DLIB_USE_CUDA DLIB_CASSERT(false, "CUDA NOT ENABLED"); +#else + DLIB_CASSERT(cuda::use_cuda(), "CUDA disabled"); #endif - DLIB_CASSERT(device_id() >= 0, "This data is host only"); copy_to_device(); device_in_use = true; return data_device.get(); @@ -129,8 +128,9 @@ namespace dlib { #ifndef DLIB_USE_CUDA DLIB_CASSERT(false, "CUDA NOT ENABLED"); +#else + DLIB_CASSERT(cuda::use_cuda(), "CUDA disabled"); #endif - DLIB_CASSERT(device_id() >= 0, "This data is host only"); copy_to_device(); host_current = false; device_in_use = true; @@ -141,8 +141,9 @@ namespace dlib { #ifndef DLIB_USE_CUDA DLIB_CASSERT(false, "CUDA NOT ENABLED"); +#else + DLIB_CASSERT(cuda::use_cuda(), "CUDA disabled"); #endif - DLIB_CASSERT(device_id() >= 0, "This data is host only"); wait_for_transfer_to_finish(); host_current = false; device_current = true; @@ -154,7 +155,14 @@ namespace dlib ) const { return host_current; } bool device_ready ( - ) const { return device_current && !have_active_transfer && device_id() >= 0; } + ) const + { +#ifdef DLIB_USE_CUDA + if (!cuda::use_cuda() && size() != 0) + return false; +#endif + return device_current && !have_active_transfer; + } size_t size() const { return data_size; } @@ -276,4 +284,3 @@ namespace dlib } #endif // DLIB_GPU_DaTA_H_ - diff --git a/dlib/cuda/gpu_data_abstract.h b/dlib/cuda/gpu_data_abstract.h index f2423dee13..b3beec6e9f 100644 --- a/dlib/cuda/gpu_data_abstract.h +++ b/dlib/cuda/gpu_data_abstract.h @@ -28,7 +28,10 @@ namespace dlib to the host do not happen before the relevant computations have completed. If DLIB_USE_CUDA is not #defined then this object will not use CUDA at all. - Instead, it will simply store one host side memory block of floats. + Instead, it will simply store one host side memory block of floats. + Similarly, if DLIB_USE_CUDA is #defined but cuda::use_cuda() == false, + then this object will be host only and will not allocate a CUDA device + memory block. THREAD SAFETY Instances of this object are not thread-safe. So don't touch one from @@ -99,10 +102,13 @@ namespace dlib ) const; /*! ensures - - returns true if and only if the device's copy of the data is current. + - returns true if and only if the device's copy of the data exists and is current. The device's data is current if there aren't any modifications to the data which were made on the host side that have yet to be copied to the device. + - if (DLIB_USE_CUDA is defined && cuda::use_cuda() == false && + size() != 0) then + - returns false. !*/ const float* host( @@ -153,6 +159,7 @@ namespace dlib /*! requires - DLIB_USE_CUDA is #defined + - cuda::use_cuda() == true ensures - returns a pointer to the device memory block of size() contiguous float values or nullptr if size()==0. @@ -167,6 +174,7 @@ namespace dlib /*! requires - DLIB_USE_CUDA is #defined + - cuda::use_cuda() == true ensures - returns a pointer to the device memory block of size() contiguous float values or nullptr if size()==0. @@ -182,6 +190,7 @@ namespace dlib /*! requires - DLIB_USE_CUDA is #defined + - cuda::use_cuda() == true ensures - This function returns the same pointer as device(), except that it never performs a host to device memory copy. Instead, it immediately marks the @@ -263,4 +272,3 @@ namespace dlib } #endif // DLIB_GPU_DaTA_ABSTRACT_H_ - diff --git a/dlib/cuda/tensor.h b/dlib/cuda/tensor.h index aaae4e836e..1f8b10e286 100644 --- a/dlib/cuda/tensor.h +++ b/dlib/cuda/tensor.h @@ -404,7 +404,10 @@ namespace dlib if ((long long)data_instance.size() < m_size) data_instance.set_size(m_size); #ifdef DLIB_USE_CUDA - cudnn_descriptor.set_size(m_n,m_k,m_nr,m_nc); + if (cuda::use_cuda()) + cudnn_descriptor.set_size(m_n,m_k,m_nr,m_nc); + else + cudnn_descriptor.set_size(0,0,0,0); #endif } @@ -640,7 +643,7 @@ namespace dlib "t.size(): "<(); inst.cudnn_descriptor->set_size(inst.m_n, inst.m_k, inst.m_nr, inst.m_nc); diff --git a/dlib/cuda/tensor_abstract.h b/dlib/cuda/tensor_abstract.h index 62f649391e..a9d56fa382 100644 --- a/dlib/cuda/tensor_abstract.h +++ b/dlib/cuda/tensor_abstract.h @@ -28,7 +28,10 @@ namespace dlib to the host do not happen before the relevant computations have completed. If DLIB_USE_CUDA is not #defined then this object will not use CUDA at all. - Instead, it will simply store one host side memory block of floats. + Instead, it will simply store one host side memory block of floats. + Similarly, if DLIB_USE_CUDA is #defined but cuda::use_cuda() == false, + then this object will be host only and will not allocate a CUDA device + memory block. Finally, the convention in dlib code is to interpret the tensor as a set of num_samples() 3D arrays, each of dimension k() by nr() by nc(). Also, @@ -151,6 +154,7 @@ namespace dlib /*! requires - DLIB_USE_CUDA is #defined + - cuda::use_cuda() == true ensures - returns a pointer to the device memory block of size() contiguous float values or nullptr if size()==0. @@ -164,6 +168,7 @@ namespace dlib /*! requires - DLIB_USE_CUDA is #defined + - cuda::use_cuda() == true ensures - returns a pointer to the device memory block of size() contiguous float values or nullptr if size()==0. @@ -179,6 +184,7 @@ namespace dlib /*! requires - DLIB_USE_CUDA is #defined + - cuda::use_cuda() == true ensures - This function returns the same pointer as device(), except that it never performs a host to device memory copy. Instead, it immediately marks the @@ -731,5 +737,3 @@ namespace dlib } #endif // DLIB_DNn_TENSOR_ABSTRACT_H_ - - diff --git a/dlib/cuda/tensor_tools.cpp b/dlib/cuda/tensor_tools.cpp index be01d38533..bc0d903e40 100644 --- a/dlib/cuda/tensor_tools.cpp +++ b/dlib/cuda/tensor_tools.cpp @@ -69,7 +69,7 @@ namespace dlib { namespace tt ) IF_DLIB_NOT_USE_CUDA( - out = sum_cols(pointwise_multiply(mat(lhs), mat(rhs))); + out = sum_cols(pointwise_multiply(mat(lhs), mat(rhs))); ) } @@ -86,9 +86,9 @@ namespace dlib { namespace tt IF_DLIB_NOT_USE_CUDA( if (add_to) - out += sum_cols(pointwise_multiply(mat(lhs), mat(rhs))); + out += sum_cols(pointwise_multiply(mat(lhs), mat(rhs))); else - out = sum_cols(pointwise_multiply(mat(lhs), mat(rhs))); + out = sum_cols(pointwise_multiply(mat(lhs), mat(rhs))); ) } @@ -351,7 +351,7 @@ namespace dlib { namespace tt ) IF_DLIB_NOT_USE_CUDA( - for (auto& x : data) + for (auto& x : data) x = cpu_impl.get_random_gaussian()*stddev + mean; ) } @@ -366,7 +366,7 @@ namespace dlib { namespace tt ) IF_DLIB_NOT_USE_CUDA( - for (auto& x : data) + for (auto& x : data) x = cpu_impl.get_random_float(); ) } diff --git a/dlib/dnn/core_abstract.h b/dlib/dnn/core_abstract.h index 8c61af5bf9..77c4ebb850 100644 --- a/dlib/dnn/core_abstract.h +++ b/dlib/dnn/core_abstract.h @@ -202,11 +202,13 @@ namespace dlib ); /*! ensures - - If dlib should use CUDA then this function returns true and false otherwise. + - If dlib will use CUDA then this function returns true and false otherwise. - This function returns true if DLIB_USE_CUDA is defined, the DLIB_DISABLE_CUDA_USE environment variable is not set to a true value, and there is an available GPU device to use. - This function always returns false if DLIB_USE_CUDA is not defined. + - A true value of DLIB_DISABLE_CUDA_USE is any value other than "", + "0", "false", "False", or "FALSE". - The value returned by this function is fixed after its first call. !*/ } @@ -2096,3 +2098,4 @@ namespace dlib } #endif // DLIB_DNn_CORE_ABSTRACT_H_ + diff --git a/dlib/test/dnn.cpp b/dlib/test/dnn.cpp index 00e299d255..1636acbfed 100644 --- a/dlib/test/dnn.cpp +++ b/dlib/test/dnn.cpp @@ -850,7 +850,7 @@ namespace IF_DLIB_USE_CUDA( input /= 2; - resizable_tensor output_cuda_a, output_cuda_b(input); + resizable_tensor output_cuda_a, output_cuda_b(input); output_cuda_a.copy_size(output_cpu_a); cuda::transpose(false, output_cuda_a, input); cuda::transpose(true, output_cuda_b, output_cuda_a); From dbc8a44d9bad7fc61245cb530c9c239a259c852a Mon Sep 17 00:00:00 2001 From: Davis King Date: Sun, 3 May 2026 19:07:02 -0400 Subject: [PATCH 36/39] more cleanup --- dlib/cuda/gpu_data.cpp | 3 +++ dlib/cuda/gpu_data_abstract.h | 2 ++ dlib/cuda/tensor_abstract.h | 2 ++ dlib/test/cublas.cpp | 2 ++ 4 files changed, 9 insertions(+) diff --git a/dlib/cuda/gpu_data.cpp b/dlib/cuda/gpu_data.cpp index 17b85288f8..3437d68cfd 100644 --- a/dlib/cuda/gpu_data.cpp +++ b/dlib/cuda/gpu_data.cpp @@ -153,6 +153,9 @@ namespace dlib void gpu_data:: async_copy_to_device() const { + if (!cuda::use_cuda()) + return; + if (!device_current) { if (device_in_use) diff --git a/dlib/cuda/gpu_data_abstract.h b/dlib/cuda/gpu_data_abstract.h index b3beec6e9f..1c16bf47e6 100644 --- a/dlib/cuda/gpu_data_abstract.h +++ b/dlib/cuda/gpu_data_abstract.h @@ -70,6 +70,8 @@ namespace dlib ); /*! ensures + - if (cuda::use_cuda() == false) then + - this function does nothing. - if (!device_ready()) then - Begins asynchronously copying host data to the device once it is safe to do so. I.e. This function will wait until any previously diff --git a/dlib/cuda/tensor_abstract.h b/dlib/cuda/tensor_abstract.h index a9d56fa382..810610e2c1 100644 --- a/dlib/cuda/tensor_abstract.h +++ b/dlib/cuda/tensor_abstract.h @@ -93,6 +93,8 @@ namespace dlib /*! ensures - This function does not block. + - if (cuda::use_cuda() == false) then + - this function does nothing. - if (the host version of the data is newer than the device's copy) then - Begins asynchronously copying host data to the device. - A call to device() that happens before the transfer completes will diff --git a/dlib/test/cublas.cpp b/dlib/test/cublas.cpp index f1fbc5b491..588ca72af6 100644 --- a/dlib/test/cublas.cpp +++ b/dlib/test/cublas.cpp @@ -58,6 +58,8 @@ namespace void perform_test ( ) { + if (!cuda::use_cuda()) return; + { cuda::cuda_data_ptr nonconst; cuda::cuda_data_ptr const_ptr(nonconst); From 77768320c6926b975b441f087d034514c71974e6 Mon Sep 17 00:00:00 2001 From: Davis King Date: Sun, 3 May 2026 19:22:12 -0400 Subject: [PATCH 37/39] cleanup --- dlib/cuda/gpu_data_abstract.h | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/dlib/cuda/gpu_data_abstract.h b/dlib/cuda/gpu_data_abstract.h index 1c16bf47e6..0c15cb5475 100644 --- a/dlib/cuda/gpu_data_abstract.h +++ b/dlib/cuda/gpu_data_abstract.h @@ -108,8 +108,7 @@ namespace dlib The device's data is current if there aren't any modifications to the data which were made on the host side that have yet to be copied to the device. - - if (DLIB_USE_CUDA is defined && cuda::use_cuda() == false && - size() != 0) then + - if (DLIB_USE_CUDA is defined && cuda::use_cuda() == false && size() != 0) then - returns false. !*/ From f743b43675a8f5c8448438f06b825235cc57b280 Mon Sep 17 00:00:00 2001 From: Davis King Date: Sun, 3 May 2026 19:34:53 -0400 Subject: [PATCH 38/39] more cleanup --- dlib/cuda/cuda_dlib.cu | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/dlib/cuda/cuda_dlib.cu b/dlib/cuda/cuda_dlib.cu index 477f917452..30c601f184 100644 --- a/dlib/cuda/cuda_dlib.cu +++ b/dlib/cuda/cuda_dlib.cu @@ -57,7 +57,7 @@ namespace dlib int get_device ( ) { - int dev = -1; + int dev = 0; if (use_cuda()) CHECK_CUDA(cudaGetDevice(&dev)); return dev; @@ -129,6 +129,9 @@ namespace dlib int peer_device_id ) : call_disable(false), device_id(device_id), peer_device_id(peer_device_id) { + if (!use_cuda()) + return; + raii_set_device set_dev(device_id); auto err = cudaDeviceEnablePeerAccess(peer_device_id, 0); @@ -3273,4 +3276,3 @@ namespace dlib } } - From 1c8455dc527a27880f74b4129ad41ef617a589df Mon Sep 17 00:00:00 2001 From: Davis King Date: Sun, 3 May 2026 19:44:03 -0400 Subject: [PATCH 39/39] more cleanup --- dlib/cuda/cuda_dlib.cu | 12 ++++++++++-- 1 file changed, 10 insertions(+), 2 deletions(-) diff --git a/dlib/cuda/cuda_dlib.cu b/dlib/cuda/cuda_dlib.cu index 30c601f184..e4f32e775d 100644 --- a/dlib/cuda/cuda_dlib.cu +++ b/dlib/cuda/cuda_dlib.cu @@ -50,8 +50,13 @@ namespace dlib int dev ) { - if (use_cuda()) - CHECK_CUDA(cudaSetDevice(dev)); + if (!use_cuda()) + { + DLIB_CASSERT(dev == 0, "dlib::cuda::set_device(id) called with an invalid device id."); + return; + } + + CHECK_CUDA(cudaSetDevice(dev)); } int get_device ( @@ -68,7 +73,10 @@ namespace dlib ) { if (!use_cuda()) + { + DLIB_CASSERT(device == 0, "dlib::cuda::get_device_name(device) called with an invalid device id."); return "CUDA_DISABLED"; + } cudaDeviceProp props; CHECK_CUDA(cudaGetDeviceProperties(&props, device));