Add cuda inference example: #155 by a-vartenkov · Pull Request #210 · KasperskyLab/knp

a-vartenkov · 2026-07-29T10:14:52Z

Adding CUDA inference example. Now CUDA backend works at least with a decent speed.

artiomn · 2026-07-31T12:54:21Z

+    if (!model_desc.model_saving_path_.empty())
+    {
+        if (!model_desc.inference_only_)
+        {
+            save_network(model_desc, network);
+        }
+        knp::framework::Network new_network = knp::framework::sonata::load_network(model_desc.model_saving_path_);
+        if (new_network.populations_count() != network.network_.populations_count()
+        || new_network.projections_count() != network.network_.projections_count())
+        {
+            std::cout << "Populations " << new_network.populations_count() << " vs. "
+                      << network.network_.populations_count() << std::endl;
+            std::cout << "Projections: " << new_network.projections_count() << " vs. "
+                      << network.network_.projections_count() << std::endl;
+        }
+        network.network_ = new_network;
+    }


Что это такое?

Сохранение и загрузка обученной модели?

artiomn · 2026-07-31T12:54:58Z

     */
    std::filesystem::path model_saving_path_;
+
+    /// A flag to not do the training.


Может, лучше в тэги?

Только если всю в теги, иначе это по-моему естественное место. А всю в теги это другая задача, по переработке примера.

artiomn · 2026-07-31T13:55:31Z

+};
+
+
+struct ValueIndex


Несколько комментов бы.

artiomn · 2026-07-31T13:58:27Z

+    gather_index_neuron_kernel<<<num_blocks, num_threads>>>(index.view(), inputs, buffer);
+    // in-place prefix sum, for each neuron the value is the number of synapses before this, starts with 0.
+    thrust::exclusive_scan(thrust::device, buffer, buffer + inputs.size_, buffer);
+    return CUDAVector<unsigned long long>{buffer, inputs.size_}; // The vector would take care of releasing


Лучше коммент наверх.

artiomn · 2026-07-31T13:59:35Z



+template <typename BaseSynapseType>
+void Network::upcast_projections()


Думаю, что в Network этого быть не должно.

А где должны быть функции, которые изменяют большой кусок Network?

Функции должны быть во фреймворке (см., там уже есть то, что работает с Network), а вот методов таких не должно быть.

artiomn · 2026-07-31T14:00:39Z

+                    is_converted = true;
+                    return core::AllProjectionsVariant{res};
+                }
+                return core::AllProjectionsVariant{proj};  // TODO: Remove unnecessary copying.


Коммент наверх.

artiomn

См. комменты.

artiomn

Продолжаем...

artiomn · 2026-08-07T13:36:12Z

+//    /**
+//     * @brief Send a message to the message bus.
+//     * @param message message to send.
+//     */
+//     template <class MessageType>
+//    __host__ __device__ void send_message(const MessageType &message);


Удалить.

artiomn · 2026-08-07T13:42:48Z

+    for (auto iter = buffer.begin(); iter != buffer.end(); ++iter)
+    {
+        LongIndex current_neuron = (*iter).first;
+        // Filling offsets for skipped neurons: if first three are missing that would be (0, 0, 0, 0, 5...


Скобку в комменте не закрыл.

artiomn · 2026-08-07T13:44:19Z

+template <class T>
+struct CUDAVectorView
+{
+    const T * const data_;
+    const unsigned long long size_;
+};
+
+template <class T>
+struct CUDAVectorMutableView
+{
+    T * const data_;
+    const unsigned long long size_;
+};
+


Потому что есть отдельно векторы, отдельно view, и если сюда будет что-то ещё добавляться, тут будет помойка.

artiomn · 2026-08-07T13:46:22Z

    {
 #if defined(__CUDA_ARCH__)
-        PRINTF_TRACE("Using has_sender on device\n");
+        // PRINTF_TRACE("Using has_sender on device\n");


FAST_ERROR_CHECK же, вроде, был?

artiomn · 2026-08-07T14:04:04Z

        size_t tail_length = end() - end_iter;
        size_t num_destruct = end() - begin_iter;
-#ifdef __CUDA_ARCH__
+    #ifdef __CUDA_ARCH__ // Device only


Вообще-то, у нас есть style guide, и никаких "вопросов читабельности" не должно возникать.

Это и нужно для того, чтобы разбивать функцию и видеть макрос, - лишняя "красота" здесь только мешает.

artiomn · 2026-08-07T14:05:20Z

+    // size_t new_heap = 128 * 1024 * 1024; // 128 Mb // TEMP!
+    // cudaDeviceSetLimit(cudaLimitMallocHeapSize, new_heap);


Тогда удали это.

artiomn · 2026-08-07T14:08:07Z



+template <typename BaseSynapseType>
+void Network::upcast_projections()


Функции должны быть во фреймворке (см., там уже есть то, что работает с Network), а вот методов таких не должно быть.

artiomn · 2026-08-11T12:05:06Z

+        }
+        knp::framework::Network new_network = knp::framework::sonata::load_network(model_desc.model_saving_path_);
+        if (new_network.populations_count() != network.network_.populations_count()
+        || new_network.projections_count() != network.network_.projections_count())


Отступ.

artiomn · 2026-08-11T12:05:27Z

@@ -45,20 +45,44 @@ template <typename Neuron>
 void run_model(const ModelDescription& model_desc)


Переписать.

artiomn · 2026-08-11T12:09:25Z

        case 0: result = extract_message_by_index<0>(msg_ptr); break;
        case 1: result = extract_message_by_index<1>(msg_ptr); break;


Сделать константами.

artiomn · 2026-08-11T12:16:30Z

+    {
+        bus.sync_with_host();
+    }
+    SPDLOG_DEBUG("Message bus 1 spikes {} impacts {}", bus.get_num_messages<cuda::SpikeMessage>(),


artiomn · 2026-08-11T12:22:11Z


 foreach (lib_name IN ITEMS "${PROJECT_NAME}" "${PROJECT_NAME}_static")
-    target_compile_options("${lib_name}" PRIVATE $<$<COMPILE_LANGUAGE:CUDA>: --expt-relaxed-constexpr -rdc=true>)
+    target_compile_options("${lib_name}" PRIVATE $<$<COMPILE_LANGUAGE:CUDA>: --expt-relaxed-constexpr -rdc=true> --generate-line-info) #Remove generate-line-info or make it debug


Убрать под DEBUG.

artiomn · 2026-08-11T12:25:50Z

+#include <knp/synapse-traits/all_traits.h>
+
+
+namespace knp::framework::projection::synapse_generators


Докстринги.

artiomn · 2026-08-11T12:27:14Z

+    /**
+     * @brief Upcast all synapses derived from SynapseType to SynapseType, leaving other synapses intact.
+     * @tparam BaseSynapseType base synapse type.
+     */
+    template <typename SynapseType>
+    void upcast_projections();


Перенести в функции.

a-vartenkov added 29 commits February 12, 2026 17:50

Fix some output typos: KasperskyLab#155

1c7c462

Change argument parsing for tests, add conversion templates: Kaspersk…

bf3c8e1

…yLab#155

Add model upcasting to example: KasperskyLab#155

867c9db

Add inference-only mode and other options: KasperskyLab#155

be3995c

Trying to merge: KasperskyLab#155

0da409c

In process of debugging

009a212

Restore tests: KasperskyLab#115

d2381ca

Fix 98s and memory overflow: KasperskyLab#155

961200a

Fixed wrong allocation for vector: KasperskyLab#155

4032f4c

Fixed population spiking: KasperskyLab#155

dd4fab8

A version of projection calculation that at least compiles without an…

d34eae5

… error: KasperskyLab#155

Fixed projection UID extraction: KasperskyLab#155

1b8df26

Fix wrong message vector initialization: KasperskyLab#155

f68f2bb

At least there are spikes and it's not segfault: KasperskyLab#155

45fbed7

Fixed Smallest Network test for CUDA backend: KasperskyLab#155

ba970c9

Fix some unremoved temporaries: KasperskyLab#155

1887093

Remove error 98: KasperskyLab#155

9d19634

Fixed synaptic impact creation: KasperskyLab#155

ed258ae

Fixed wrong indexing: KasperskyLab#155

6855c51

Fix GPU inference: KasperskyLab#155

ef77027

Parallelize populations: KasperskyLab#155

50f1526

Fix impact calculations: KasperskyLab#155

7691485

Some small optimizations for the vector: KasperskyLab#155

af8ef75

Changed the new message bus to be compilable: KasperskyLab#155

5b87183

One iteration is correct: KasperskyLab#155

12d91cf

Inference seems to be working: KasperskyLab#155

cc04a19

Change some synchronizes: KasperskyLab#155

e9ea2d9

Merging master: KasperskyLab#155

075865a

Finish merging: KasperskyLab#155

6a6b319

artiomn requested changes Jul 30, 2026

View reviewed changes

Comment thread examples/mnist-client/inference.cpp

artiomn requested changes Jul 31, 2026

View reviewed changes

a-vartenkov added 4 commits August 4, 2026 13:10

Change file names, remove _alt: KasperskyLab#155

09b2114

Rename files, fix some tabs: KasperskyLab#155

b5be361

Some files and folders were renamed: KasperskyLab#155

43327cc

Move comment, add an empty string: KasperskyLab#155

23ba7f0

artiomn requested changes Aug 5, 2026

View reviewed changes

Comment thread examples/mnist-learn/models/blifat/prepare_network_for_inference.cpp

Comment thread examples/mnist-learn/main.cpp

a-vartenkov added 3 commits August 7, 2026 13:52

Changing types and file names: KasperskyLab#155

96e438d

Printfs moved to debug: KasperskyLab#155

e9df268

Small changes to vector: KasperskyLab#155

0d171e1

artiomn changed the title ~~Feature/#155 add cuda inference example~~ Add cuda inference example: #155 Aug 7, 2026

artiomn requested changes Aug 7, 2026

View reviewed changes

Fix debugs and tabs at cuda bus: KasperskyLab#155

07ff4c8

artiomn requested changes Aug 11, 2026

View reviewed changes



		template <typename BaseSynapseType>
		void Network::upcast_projections()

		// size_t new_heap = 128 * 1024 * 1024; // 128 Mb // TEMP!
		// cudaDeviceSetLimit(cudaLimitMallocHeapSize, new_heap);

		@@ -45,20 +45,44 @@ template <typename Neuron>
		void run_model(const ModelDescription& model_desc)

		case 0: result = extract_message_by_index<0>(msg_ptr); break;
		case 1: result = extract_message_by_index<1>(msg_ptr); break;

		#include <knp/synapse-traits/all_traits.h>


		namespace knp::framework::projection::synapse_generators

		};


		struct ValueIndex

Uh oh!

Conversation

a-vartenkov commented Jul 29, 2026

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

artiomn left a comment

Choose a reason for hiding this comment

Uh oh!

artiomn left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

artiomn Aug 7, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

artiomn Aug 7, 2026 •

edited

Loading