Неразмеченные K-Means с использованием Python Numba и CUDA C

Улучшение с использованием Python Numba и CUDA C неразмеченные алгоритмы K-Means

Как ускорить анализ данных на 1600 раз по сравнению с Scikit-Learn (с кодом!)

Изображение, созданное с помощью Midjourney на основе рисунка автора

Параллельное выполнение рабочих нагрузок анализа данных может быть сложной задачей, особенно когда нет эффективной готовой реализации для вашего конкретного случая использования. В этом руководстве я расскажу о принципах написания ядер CUDA на C и Python Numba, а также о том, как эти принципы могут быть применены к классическому алгоритму кластеризации k-средних. По окончании этой статьи вы сможете написать настраиваемую параллельную реализацию пакетной кластеризации k-средних как на C, так и на Python, достигая ускорения до 1600 раз по сравнению с обычной реализацией scikit-learn. Если вы хотите сразу перейти к коду, он доступен в Colab.

Введение

Мы изучим две платформы для параллельного выполнения задач на графических процессорах NVIDIA: Python Numba и CUDA C. Затем мы реализуем алгоритм пакетной кластеризации k-средних на обоих платформах и проведем их сравнение с scikit-learn.

Numba — это более простая система обучения для тех, кто предпочитает Python перед C. Numba компилирует части вашего кода в специализированные функции CUDA, называемые ядрами. CUDA C предоставляет более тонкую настройку и находится на более низком уровне абстракции. Примеры в Colab разработаны таким образом, чтобы тесно соответствовать реализациям алгоритмов, поэтому после понимания одного из них вы сможете легко разобраться с другим.

Это руководство выросло из нашей собственной настраиваемой реализации пакетной кластеризации k-средних, для которой нам понадобилось кластеризовать миллионы небольших наборов данных параллельно, и мы не смогли найти готовую библиотеку для этого.

Эта статья разделена на четыре основные секции: CUDA Basics (Основы CUDA), Parallelized K-means Algorithm (Параллельный алгоритм k-средних), Python Numba и CUDA C. Я постараюсь делать материал относительно самодостаточным и кратко повторю концепции по мере необходимости.

Основы CUDA

Центральные процессоры (CPUs) разработаны для быстрой последовательной обработки, а графические процессоры (GPUs) — для массовой параллельности. Для нашего алгоритма нам нужно было запустить k-средние на миллионах маленьких независимых наборов данных, что идеально подходит для реализации на GPU.

Наша реализация будет использовать CUDA (Compute Unified Device Architecture) — библиотеку на языке C и платформу вычислений, разработанные NVIDIA для использования возможностей GPU для параллельных вычислений.

Для понимания ядер GPU и устройства функций, приведем некоторые определения:

Поток — это единичная последовательность инструкций, которую можно выполнить независимо.

Ядро — это вычислительная единица, которая может выполнить один поток.

Виток — это самая маленькая единица планирования потоков. Каждый виток состоит из 32 потоков и размещает их на ядрах.

Мультипроцессор (также называемый потоковым мультипроцессором, SM) состоит из определенного количества ядер.

Ядро GPU — это функция, которая предназначена для параллельного выполнения множества потоков GPU. Ядра функций вызываются процессором, а выполняются на GPU.

Функция устройства — это функция, которую может вызывать ядро функции или другая функция устройства. Эта функция выполняется на GPU.

Эти ядра организованы в иерархию сетки, блоков и потоков. В контексте GPU каждый поток связан с одним ядром и выполняет копию ядра.

Структура данных GPUFunction[(x, y), z]. Изображение автора.

Блок – это коллекция потоков, которая работает на одном мультипроцессоре.

Сетка – это абстрактный массив для организации блоков потоков. Сетки используются для сопоставления экземпляров ядра с потоками по индексу.

На GPU есть два типа памяти: глобальная память и разделяемая память.

Глобальная память хранится в DRAM (динамическая оперативная память). Все потоки могут получить доступ к ее содержимому, но это связано с большой задержкой памяти.

Разделяемая память хранится в кэше и является приватной для потоков внутри одного блока. К ней можно получить доступ гораздо быстрее, чем к глобальной памяти.

Параллельный алгоритм K-средних

Теперь давайте познакомимся с алгоритмом K-средних. K-средние – это безнадзорный алгоритм кластеризации, который разбивает набор данных на k различных, не перекрывающихся кластеров. Учитывая набор точек данных, мы сначала инициализируем k центроидов, или начальные центры кластеров:

Инициализация центроидов (k=3). Изображение от автора.

Затем, после выбора начальных центроидов, последовательно выполняем два шага:

  1. Шаг назначения: Каждая точка данных назначается ближайшему центроиду на основе евклидова расстояния.
  2. Шаг обновления: Положение центроидов переназначается как среднее из всех точек, назначенных этому центроиду на предыдущем шаге.

Эти шаги повторяются до сходимости, или когда положение центроидов больше не изменяется значительно. В качестве результата получаем набор координат центроидов кластеров k вместе с массивом, помечающим индексы кластеров (называемый метками) для каждой из исходных точек данных.

Алгоритм K-средних (k = 3). Изображение от автора.

Для большого набора данных инициализация центроидов может значительно влиять на результат работы алгоритма. Поэтому программа пробует несколько начальных центроидов, называемых начальными зернами, и возвращает результат от лучшего зерна. Каждое зерно выбирается из начального набора данных без повторений, то есть ни один начальный центроид не будет повторяться. Оптимальное количество зерен для нашего алгоритма составляет треть от количества точек данных. В нашей программе, поскольку мы запускаем K-средних на отдельных строках из 100 точек данных, оптимальное количество зерен будет равно 33.

В нашей функции K-средних блоки представляются одним миллионом строк, а потоки представляют зерна. Потоки в блоке организованы в волны, которые являются наименьшей единицей планирования потоков в аппаратной архитектуре. Каждая волна состоит из 32 потоков, поэтому оптимально устанавливать размеры блоков как кратные 32. Затем каждый блок выводит данные от зерна, у которого наименьшая инерция, измеряемая суммой евклидового расстояния между центрами кластеров и их присвоенными точками.

Параллельный K-средних - сторона GPU. Изображение от автора.

Здесь ссылка на Colab, где вы можете следовать вместе с помощью Python или C.

глобальные переменные: numRows, lineSize, numClustersdef hostKMeans:    inputData = initializeInputData(numRows, lineSize)    outputCentroids = createEmptyArray(numRows, numClusters)    outputLabels = createEmptyArray(numRows, lineSize)        sendToDevice(inputData, outputCentroids, outputLabels)    cuda_kmeans[blockDimensions, threadDimensions](inputData, outputCentroids, outputLabels)    waitForKernelCompletion()    copyFromDevice(outputCentroids, outputLabels)

В нашем алгоритме K-средних мы начинаем с установки глобальных переменных. Нам потребуется обращаться к ним как с ЦП, так и с ГП.

В то время как глобальные переменные доступны из ядра, ядро не может напрямую возвращать значения хосту. Чтобы обойти это ограничение, часть нашего кода на ЦП передает два пустых массива ядру вместе с входными данными. Эти массивы будут использоваться для копирования окончательного массива центров и меток обратно на ЦП.

После создания наших структур данных мы вызываем ядро, задавая размеры сетки и блока в виде кортежа. Вызов ядра включает передачу наших данных, выделение памяти для центроидов и меток, а также состояния для начальной случайной инициализации центроидов.

def KMeansKernel(data, outputCentroids, outputLabels)    row = currentBlock()    seed = currentThread()    sharedInputRow = sharedArray(shape=(lineSize))    sharedInertia = sharedArray(shape=(numSeeds))    sharedCentroids = sharedArray(shape=(numSeeds, numClusters))    sharedLabels = sharedArray(shape=(numSeeds, lineSize))    sharedInputRow = data[row]    synchronizeThreads()    if seed == 0        centroids = initializeCentroids(data)    synchronizeThreads()        KMeansAlgorithm(sharedInputRow, sharedCentroids, sharedLabels)        sharedInertia[Seed] = calculateInertia(sharedInputRow, sharedCentroids, sharedLabels)        synchronizeThreads()    if seed == 0        minInertiaIndex = findMin(Inertia)    sharedOutputCentroids = centroids[minInertiaIndex]    sharedOutputLabels = labels[minInertiaIndex]

Около устройства (или ГПУ) наш код существует одновременно на всей сетке. Чтобы определить, где мы находимся в сетке, мы обращаемся к индексам блока и потока.

На ГПУ память по умолчанию является глобальной памятью, которая хранится на DRAM. Общая память является личной для потоков в пределах одного блока.

Передачей одной строки данных (которую все потоки в блоке должны прочитать) в разделяемую память мы сокращаем время доступа к памяти. В функции cuda_kmeans мы создаем разделяемую память для хранения строки данных, центроидов, меток и меры точности для каждого семени, называемой инерцией.

В нашей программе каждый поток соответствует одному семени, и все потоки в блоке работают с одной строкой данных. Для каждого блока один поток последовательно создает 32 семени и агрегирует их результаты в одну структуру данных для других потоков в блоке.

Когда последующий шаг алгоритма полагается на завершение этой агрегации, потокам необходима синхронизация с использованием встроенной функции CUDA syncthreads(). NB: нужно быть очень осторожными с расположением вызовов syncthreads(), так как попытка синхронизации потоков, когда все они еще не успели завершиться, может привести к взаимоблокировкам и зависанию всей программы.

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

def KMeansDevice(dataRow, centroids, labels)    seed = currentThread()    centroidsRow = centroids[seed]    labelsRow = labels[seed]     centroidsRow = sort(centroidsRow)    yardStick = computeYardstick(sortedCentroids)     oldCentroids = localArray(shape=(numSeeds, numClusters))    for iteration in range(100):        if converged(oldCentroids, centroidsRow)            break        oldCentroids = copy(centroidsRow)        assignLabels(dataRow, centroidsRow, labelsRow)        updateCentroids(dataRow, centroidsRow, labelsRow)

Из функции cuda_kmeans мы затем вызываем фактический алгоритм k-средних, передавая вновь созданную разделяемую память. В нашем алгоритме k-средних мы начинаем с выбора начальных центроидов и их сортировки по возрастанию. Мы итеративно присваиваем точки данных ближайшему центроиду и обновляем их расположение до сходимости.

Чтобы определить, достигнута ли сходимость, мы используем вспомогательную функцию find_yard_stick. Эта функция вычисляет и возвращает наименьшее расстояние между двумя начальными центроидами (yard_stick). Наше условие сходимости выполнено, когда ни один из центроидов не перемещается на расстояние более yard_stick раз epsilon в течение одной итерации.

После достижения сходимости мы возвращаемся к функции cuda_kmeans. Здесь мы определяем лучшее семя, вычисляя квадратичное евклидово расстояние между каждым центроидом и его точками данных. Семя с наиболее эффективной группировкой, указанной наименьшей инерцией, считается лучшим. Затем мы берем центроиды и метки из этого семени и копируем их в одну строку в наши выходные массивы. После завершения всех блоков эти выводы затем копируются обратно на хост (ЦП).

Трансферы данных в алгоритме K-средних. Изображение автора.

Введение в Numba

Наиболее простым способом создания пользовательских ядер является использование Numba. Numba – это библиотека Python, которая может использоваться для компиляции Python-кода в ядра CUDA.

Уровни абстракции. Изображение автора.

Под поверхностью Numba взаимодействует с CUDA. Для параллелизации вашего кода Numba компилирует ваш код, предназначенный для GPU, в ядро и передает его на GPU, разделяя логику программы на две основные части:

  1. Код уровня ЦП
  2. Код уровня ГП

Используя Numba, вы отделяете и передаете последовательные и параллелизуемые части кода ЦП и ГП. Чтобы скомпилировать функцию для GPU, программист использует декоратор @cuda.jit над определением функции, тем самым превращая эту функцию в ядро, которое вызывается с ЦП (хоста), но выполняется параллельно на ГП (устройстве).

Python Numba

Ссылка на Colab.

Numba служит мостом между кодом Python и платформой CUDA. Поскольку код Python практически идентичен псевдокоду алгоритма выше, я приведу всего несколько примеров ключевого синтаксиса.

cuda_kmeans[(NUM_ROWS,), (NUM_SEEDS,)](input_rows, output_labels, output_centroids, random_states)

После создания необходимых глобальных переменных и структур данных мы можем вызвать ядро cuda_kmeans с хоста. Numba требует два кортежа для размерности блоков и потоков. Поскольку мы будем использовать одномерные блоки и потоки, второй индекс в каждом кортеже пуст. Мы также передаем наши структуры данных и массив случайных состояний для инициализации случайного семени.

@cuda.jit()def cuda_kmeans(input, output_labels, output_centroids, random_states):    row = cuda.blockIdx.x    seed = cuda.threadIdx.x    shared_input_row = cuda.shared.array(shape=(LINE_SIZE), dtype=np.float32)    shared_inertia = cuda.shared.array(shape=(NUM_SEEDS), dtype=np.float32)    shared_centroids = cuda.shared.array(shape=(NUM_SEEDS, NUM_CLUSTERS), dtype=np.float32)    shared_labels = cuda.shared.array(shape=(NUM_SEEDS, LINE_SIZE), dtype=np.int32)    if seed == 0:        get_initial_centroids(shared_input_row, shared_centroids, random_states)    cuda.syncthreads()    ...    kmeans(shared_input_row, shared_labels, shared_centroids)

Мы используем декоратор Numba @cuda.jit() для отметки функции для компиляции для ГП. Используя обозначения cuda.blockIdx.x и cuda.threadIdx.x, мы получаем текущий индекс ядра. Общие массивы создаются с помощью cuda.shared.array с двумя аргументами: формой и типом, которые должны быть известны во время компиляции. После получения центроидов и заполнения строки данными мы вызываем функцию kmeans, заполняем общие массивы и делаем вызов cuda.syncthreads().

@cuda.jit(device=True)def kmeans(data_row, output_labels, output_centroids):     seed = cuda.threadIdx.x    labels_row = output_labels[seed]    centroids_row = output_centroids[seed]        ...    old_centroids = cuda.local.array(shape=(NUM_CLUSTERS), dtype=np.float32)    for iteration in range(NUM_ITERATIONS):            if iteration > 0:                if converged(centroids_row, old_centroids, yard_stick * EPSILON_PERCENT, iteration):                    break      # Назначить метки и обновить центроиды

K-средних – это вспомогательная функция, так как она вызывается из функции ядра. Поэтому мы должны указать device=True в декораторе: @cuda.jit(device=True). Функция k-средних затем получает текущую строку для меток и центроидов и выполняется до сходимости.

С помощью всего лишь десятка строк кода и немного усилий, ваш код на Python может стать оптимизированным ядром, готовым для параллельного использования.

Наш параллельный k-means значительно сокращает время вычислений, но обертывание и компиляция языка Python, как правило, не являются оптимальными. Чтобы узнать, ускорило ли написание кода на C наш проект, я погрузился в мир CUDA C.

Введение в C

В Python выделение памяти и определение типов происходит автоматически, тогда как в C память может выделяться в стеке или куче, и для этого требуется явное определение типа и выделение фиксированного объема памяти. Память в стеке автоматически выделяется и освобождается компилятором, тогда как память в куче выделяется вручную во время выполнения с использованием функций, таких как malloc(), и программист несет ответственность за ее освобождение.

Указатели – это инструменты, которые содержат адреса в памяти переменных. Тип данных, на который указывает указатель, определяется при объявлении. Указатели указываются с помощью знака «*». Чтобы получить адрес переменной, называемый приемом ссылки, используется амперсанд («&»). Чтобы получить значение из указателя, вы снова разыменовываете указатель с помощью звездочки.

Двойной указатель, или указатель на указатель, хранит адрес указателя. Это полезно для изменения адресов других указателей при передаче массивов в функции. При передаче массивов в функции они передаются как указатели на их первый элемент, без информации о размере, полагаясь на арифметику указателей для индексирования через массив. Чтобы вернуть массив из функции, нужно передать адрес указателя с помощью & и принять его с двойным указателем **, что позволяет изменить адрес исходного указателя, таким образом, передавая массив.

int var = 100; // объявление типaint *ptr = &var; // использование указателя и ссылкиint **double_ptr = &ptr; // пример двойного указателяprintf(“Двойная десериализация double_ptr и ptr: %d %d \n:”, **double_ptr, *ptr)int *ptr = 100; // инициализация указателя типа int

CUDA C

Ссылка на Colab.

CUDA – вычислительная платформа, которая использует графические процессоры NVIDIA для параллелизации сложных вычислительных задач. На основе вашего вновь обретенного владения C (шутка), структура кода CUDA C буквально идентична структуре псевдокода, которую мы рассмотрели.

На стороне процессора мы устанавливаем несколько констант, которые сообщают алгоритму ожидания, импортируем библиотеки, инициализируем переменные и определяем некоторые макросы.

#define NUM_ROWS 00000        // размерность y для нашего набора данных, количество блоков#define LINE_SIZE 100         // размерность x для нашего набора данных#define NUM_ITERATIONS 100    // максимальное количество итераций#define NUM_CLUSTERS 3        // Мы запускаем k = 3#define MAX_INPUT_VALUE 100   // Верхняя граница данных#define NUM_SEEDS 32          // Количество центроидов/потоков составляет 1/3 от LINE_SIZE#define EPSILON_PERCENT 0.02  // Условие сходимостиvoid initInputData(float** input) {    srand(1);     // выделить память для данных        ... // инициализировать данные, используя malloc и rand    // выделить память на GPU    cudaMalloc(input, NUM_ROWS * LINE_SIZE * sizeof(float));     // скопировать данные с памяти CPU (sample_data) в память GPU    cudaMemcpy(*input, sample_data, NUM_ROWS * LINE_SIZE * sizeof(float), cudaMemcpyHostToDevice);    free(sample_data);}int main() {    float* inputData; // инициализировать входные данные, размерности будут NUM_ROWS x LINE SIZE    initInputData(&inputData); // разыменовать и передать в функцию    // инициализировать выходные метки и центроиды    cudaExtent labelsExtent = make_cudaExtent(sizeof(int), LINE_SIZE, NUM_ROWS);    cudaPitchedPtr outputLabels; // создать указатель, необходимый для следующего вызова    cudaMalloc3D(&outputLabels, labelsExtent); // выделить память на GPU        cudaExtent centroidsExtent = make_cudaExtent(sizeof(float), NUM_CLUSTERS, NUM_ROWS);    cudaPitchedPtr outputCentroids; // создать указатель, необходимый для следующего вызова    cudaMalloc3D(&outputCentroids, centroidsExtent); // выделить память на GPU    cuda_kmeans <<<NUM_ROWS, NUM_SEEDS>>> (inputData, outputLabels, outputCentroids);    cudaDeviceSynchronize();        ... // скопировать вывод со счетчика обратно на хост}

Давайте разберем различия.

Основная функция начинается с инициализации данных путем создания указателя и передачи его в initInputData как адреса указателя. Функция получает inputData как указатель на указатель (float** input), что позволяет функции изменять адрес, хранящийся в исходном указателе. Затем ввод указывается на адрес памяти GPU, который инициализируется с использованием cudaMalloc и заполняется с использованием cudaMemcpy, копируя данные из временного массива хоста sample_data, уже заполненного случайными числами.

Затем код выделяет память на устройстве для хранения результатов от функции к-средних. Функция использует make_cudaExtent для создания объекта cudaExtent, который предназначен для инкапсуляции размеров многомерного массива.

Тип cudaPitchedPointer используется для определения указателя, который может адресовать эту установленную пространство памяти. Этот тип указателя специально разработан для работы с памятью, выделенной с помощью cudaMalloc3D, который принимает как cudaPitchedPtr, так и объект cudaExtent, чтобы выделить линейную память на GPU.

cuda_kmeans <<<NUM_ROWS, NUM_SEEDS>>> (inputData, outputLabels, outputCentroids);

Входя в код GPU, мы определяем сетку так, чтобы каждый блок соответствовал строке данных, а каждый поток – сиду.

Семена инициализируются одним потоком в каждом блоке, обеспечивая полностью разные семена.

__global__ void cuda_kmeans(float* input, cudaPitchedPtr outputLabels, cudaPitchedPtr outputCentroids) {    int row = blockIdx.x;    int seed = threadIdx.x;    // shared memory is shared between threads in blocks    __shared__ float input_shm[LINE_SIZE];    __shared__ float error_shm[NUM_SEEDS];    __shared__ float seed_centroids[NUM_SEEDS][NUM_CLUSTERS];    __shared__ int seed_labels[NUM_SEEDS][LINE_SIZE];        ... // get a single row of data    ... // populate input_shm    ... // populating the struct core_params    // the actual k-means function    kmeans(core_params);     // find seed with smallest error    calcError(core_params);    __syncthreads();    if (seed == 0) {        int* labels_line = LABELS_LINE(outputLabels, row);        float* centroids_line = CENTROIDS_LINE(outputCentroids, row);        labels_line[threadIdx.x] = seed_labels[seed][threadIdx.x];        centroids_line[threadIdx.x] = seed_centroids[seed][threadIdx.x];    }}

Код CUDA C использует разделяемую память для данных, центроидов, меток и ошибок. Однако, в отличие от Python, код принимает указатели на разделяемую память и сохраняет их в структуре, которая является просто методом передачи переменных массово. Наконец, cuda_kmeans вызывает фактический алгоритм к-средних и передает core_params.

__device__ void kmeans(core_params_t& core_params) {    DECLARE_CORE_PARAMS(core_params);    getInitialCentroids(core_params);    sort_centroids(centroids, num_clusters);    float yard_stick = findYardStick(core_params);    float* oldCentroids = (float*)malloc(NUM_CLUSTERS * sizeof(float));    struct work_params_t work_params;    work_params.min = find_min(line, LINE_SIZE);    work_params.max = find_max(line, LINE_SIZE);    work_params.aux_buf1 = (int*)malloc(NUM_CLUSTERS * sizeof(int));    work_params.aux_buf2 = (int*)malloc(NUM_CLUSTERS * sizeof(int));    work_params.aux_buf3 = (float*)malloc(NUM_CLUSTERS * sizeof(float));    for (int iterations = 0; true; iterations++) {        bool stop = (iterations > 100) || (iterations > 0 && (converged(core_params, oldCentroids, yard_stick * EPSILON_PERCENT)));        if (stop)            break;        memcpy(oldCentroids, core_params.centroids, NUM_CLUSTERS * sizeof(float));        getLabels(core_params);        getCentroids(core_params, work_params);    }    free(work_params.aux_buf1);    free(work_params.aux_buf2);    free(work_params.aux_buf3);    free(oldCentroids);}

В функции устройства, перед чем-либо, мы извлекаем значения из структуры core_params в переменные с помощью макроса DECLARE_CORE_PARAMS.

Затем мы выполняем ту же алгоритм к-средних, что и в Python, с единственной разницей – мы передаем структуру вместо переменных и должны управлять памятью, указателями и типами.

Бенчмарки

Для сравнения наших алгоритмов с непараллельным k-средних, мы импортируем модуль scikit-learn k-средних.

Для нашего бенчмарка мы запускаем 100 000 строк по 100 столбцов с тремя кластерами. Поскольку в scikit-learn нет параллельного k-средних для разных строк, мы запускаем строки последовательно в цикле for.

Для наших бенчмарков в colab мы используем бесплатный экземпляр T4 GPU Colab.

Изображение от автора.

Результаты хорошие – код Python Numba в два раза быстрее непараллельного кода CPU, а код CUDA C в три раза быстрее. Ядерные функции легко масштабируются и алгоритм можно изменять для поддержки кластеризации в более высоких измерениях.

Обратите внимание, что случайно генерируемая начальная генерация центроидов как в C, так и в Python, не полностью оптимизирована для использования всех ядер. При улучшении алгоритма Python может примерно так же равняться времени выполнения кода C.

Время работы на бесплатном Colab T4 GPU 23.11.2023 года. Изображение от автора.

После выполнения функции k-средних сто раз на разных данных и записи производительности – мы замечаем, что первая итерация значительно медленнее из-за времени, необходимого для компиляции как C, так и Python в Colab.

Заключение

Теперь вы должны быть готовы писать свои собственные кастомные ядра GPU! Одним из оставшихся вопросов может быть – следует ли использовать CUDA C или Numba для параллельной обработки данных? Все зависит. Оба намного быстрее, чем стандартный scikit-learn. Несмотря на то, что в моем случае пакетная реализация k-средних на CUDA C оказалась примерно в 3,5 раза быстрее, чем эквивалентная реализация с использованием Numba, Python предлагает некоторые важные преимущества, такие как читаемость и меньшая зависимость от специализированных навыков программирования на C в командах, которые в основном работают на Python. Кроме того, время выполнения вашей конкретной реализации будет зависеть от того, насколько хорошо вы оптимизировали свой код, чтобы, например, не вызывать сериализованные операции на GPU. В заключение, если вы новичок как в C, так и в параллельном программировании, я рекомендую начать с Numba для прототипирования вашего алгоритма, а затем перевести его в CUDA C, если вам требуется дополнительное ускорение.

Ссылки

  1. Scikit-learn: Обучение машин на языке Python, Педрегоза и др., JMLR 12, pp. 2825–2830, 2011.
  2. NVIDIA, Vingelmann, P. & Fitzek, F.H.P., 2020. CUDA, релиз: 10.2. 89, Доступно на: https://developer.nvidia.com/cuda-toolkit.
  3. Lam, Siu Kwan, Antoine Pitrou, и Stanley Seibert. «Numba: Компилятор Python на основе LLVM». Материалы Второго семинара по инфраструктуре компилятора LLVM в вычислительных системах высокой производительности. 2015.
  4. Harris, C.R., Millman, K.J., van der Walt, S.J. и др. «Массивное программирование с NumPy.» Nature 585, 357–362 (2020). DOI: 10.1038/s41586–020–2649–2. (Ссылка на издателя).