Как разделить входную матрицу на подматрицы для параллельной обработки графическим процессором?

Question

Я новичок в stackoverflow и я рад опубликовать свой первый вопрос.Я на самом деле новичок в мире программирования, но в любом случае у меня есть некоторые базовые знания, я на самом деле использую Python для решения проблемы классификации.У меня большой набор данных (в моей реальной задаче n = 60326), и каждый вектор имеет около 416 измерений.Мне нужно рассчитать манхэттенское расстояние между этими различными векторами, чтобы классифицировать мой набор данных на основе сходства (взять случайный опорный вектор и объединить самые близкие векторы с ним, где диапазон расстояний между 0 и 1) я уже знаком с Kmeans и базовой кластеризацией MLалгоритмы ... Моя настоящая проблема заключается в том, что мне нужно ускорить время, используя графический процессор (CUDA) для первоначального вычисления матрицы расстояний до начала классификации, которая имеет размер n² (60326 x 60326, мы можем уменьшить ее до n² / 2потому что это симметричная матрица) Поэтому моя актуальная проблема заключается в том, как реализовать CUDA в этом случае, я уже установил пакет CUDA с ANACONDA.

Я начал с параллельной обработки ЦП, и она вызвала ошибку памяти

import pandas as pd
import time  
import time
import numpy as np
import random
from scipy.spatial import distance
from sklearn.metrics.pairwise import pairwise_distances
signatures=pd.read_csv("C:\\Users\\YMH1\\Documents\\Reduce Java code\\BOBST.txt", sep=' ',header=None,usecols=[*range(2,417)])
PartName = pd.read_csv('C:\\Users\\YMH1\\Documents\\Reduce Java code\\misumi_new.txt', sep=' ', header=None, usecols=[*range(0,1)])
signatures_vector=np.array(signatures)
PartName_vector=np.array(PartName)
D = pairwise_distances(X = signatures_vector, metric = 'manhattan', n_jobs = -1)
print(D)

Теперь я пытаюсь реализовать CUDA, так как это ускоряет время, поэтому я написал следующее:

from __future__ import division
from numba import cuda, float32
import pandas as pd
import math
signatures=pd.read_csv("C:\\Users\\YMH1\\Documents\\Reduce Java code\\BOBST.txt", sep=';',header=None,usecols=[*range(2,418)])
signatures_vector=np.array(signatures)
@cuda.jit
def manhattan(an_array):
  x, y = cuda.grid(2)
  "Here we define the Manhattan distance"
  return an_array
data=signatures_vector
threadsperblock = (16, 16)
blockspergrid_x = math.ceil(data.shape[0] / threadsperblock[0])
blockspergrid_y = math.ceil(data.shape[1] / threadsperblock[1])
blockspergrid = (blockspergrid_x, blockspergrid_y)
manhattan[blockspergrid, threadsperblock](data)
print(data)

Мой вопрос заключается в том, как определить манхэттенскую норму, когда мы используем CUDA, какие изменения необходимо внести в def manhattan И большое спасибо

Answer 1

Мой вопрос заключается в том, как определить манхэттенскую норму при использовании CUDA, какие изменения необходимо внести в def manhattan

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

Одна из проблем, с которой вы, вероятно, столкнетесь, - это пространство памяти.Если мы предположим, что вывод метрики расстояния (т.е. элементов матрицы) выражается как обычная величина питона, это, вероятно, займет 8 байтов памяти.Для указанного размера (60326) это означает, что матрица будет занимать 60326 * 60326 * 8 байт, что составляет почти 30 ГБ.Даже если мы предположим, что вы храните только половину этого, и даже если мы примем 32-битную сумму абсолютных разностей, это все равно будет более 7 ГБ необходимого хранилища.

Когда я попытался запустить такой тест с использованием метода sckit-learn, у меня возникли проблемы даже на машине с 128 ГБ системной памяти:

# cat t5.py
import numpy as np
import random
from scipy.spatial import distance
from sklearn.metrics.pairwise import pairwise_distances

vector_length = 416
num_signatures = 60000
sig_pattern = np.array([0,1,2,3], dtype=np.float32).reshape(4,1)
signatures = np.tile(sig_pattern,(num_signatures, vector_length//sig_pattern.shape[1]))
E = pairwise_distances(signatures, metric = 'manhattan', n_jobs = -1)
print(E[:8,:8])
# time python t5.py

Traceback (most recent call last):
  File "t5.py", line 17, in <module>
    E = pairwise_distances(signatures, metric = 'manhattan', n_jobs = -1)
  File "/root/anaconda2/lib/python2.7/site-packages/sklearn/metrics/pairwise.py", line 1247, in pairwise_distances
    return _parallel_pairwise(X, Y, func, n_jobs, **kwds)
  File "/root/anaconda2/lib/python2.7/site-packages/sklearn/metrics/pairwise.py", line 1096, in _parallel_pairwise
    for s in gen_even_slices(Y.shape[0], n_jobs))
  File "/root/anaconda2/lib/python2.7/site-packages/sklearn/externals/joblib/parallel.py", line 789, in __call__
    self.retrieve()
  File "/root/anaconda2/lib/python2.7/site-packages/sklearn/externals/joblib/parallel.py", line 699, in retrieve
    self._output.extend(job.get(timeout=self.timeout))
  File "/root/anaconda2/lib/python2.7/multiprocessing/pool.py", line 572, in get
    raise self._value
multiprocessing.pool.MaybeEncodingError: Error sending result:
'[array([[ 0.,  416.,  832., ...,  416.,  832., 1248.],
       [ 416.,    0.,  416., ...,    0.,  416.,  832.],
       [ 832.,  416.,    0., ...,  416.,    0.,  416.],
       ...,
       [ 416.,    0.,  416., ...,    0.,  416.,  832.],
       [ 832.,  416.,    0., ...,  416.,    0.,  416.],
       [1248.,  832.,  416., ...,  832.,  416.,    0.]])]'. Reason: 'OverflowError('cannot serialize a string larger than 2 GiB',)'

real    31m47.361s
user    405m28.155s
sys     8m19.851s

В этой ситуации этоказалось, заняло около 30 минут, чтобы вычислить на моей машине.Выходная тестовая матрица выглядела примерно корректно, но Python выдал ошибку, потому что некоторое промежуточное представление было> 2 ГБ.хорошо.

Однако выполняемая операция относительно проста.И, согласно моему тестированию, он может работать значительно быстрее, чем метод numpy / scikit-learn.

Вот рабочий пример:

# cat t4.py
import numpy as np
import numba as nb
from numba import cuda,float32

vector_length = 416
num_vecs_per_block = 8
sm_size = vector_length*num_vecs_per_block
#num_sig must be divisible by num_vecs_per_block
@cuda.jit('void(float32[:,:], float32[:,:], int32, int32)')
def manhattan(signatures, distances, num_sig, vec_len):
    sm = cuda.shared.array(sm_size, float32)
    temp = cuda.local.array(num_vecs_per_block, dtype = float32)
    bid = cuda.blockIdx.x
    tid = cuda.threadIdx.x
    bdim = cuda.blockDim.x
# load shared memory with vectors for comparison
    if tid < num_vecs_per_block:
        for i in range(vec_len):
            sm[i*num_vecs_per_block+tid] = signatures[i, bid*num_vecs_per_block+tid];
    cuda.syncthreads()
#block-stride loop through the vector array
# the addition below to tid results in only elements above the diagonal being computed
# since the output matrix is symmetric
    svec = tid +(bid*num_vecs_per_block)
    while svec < num_sig:
        for i in range(num_vecs_per_block):
            temp[i] = 0
        for i in range(vec_len):
            src = signatures[i,svec]
            for j in range(num_vecs_per_block):
                #elementwise difference
                sdist = src - sm[i*num_vecs_per_block + j]
                #absolute value
                if sdist < 0:
                    sdist *= -1
                #sum
                temp[j] += sdist
        for i in range(num_vecs_per_block):
            distances[bid*num_vecs_per_block+i, svec] = temp[i]
        svec += bdim

num_signatures = 60000
sig_pattern = np.array([0,1,2,3], dtype=np.float32)
signatures = np.tile(sig_pattern,(num_signatures//sig_pattern.shape[0], vector_length))
distances  = np.zeros((num_signatures, num_signatures), dtype=np.float32)
threadsperblock = 1024
blockspergrid   = num_signatures//num_vecs_per_block
d_signatures = cuda.to_device(signatures)
d_distances = cuda.device_array_like(distances)
manhattan[blockspergrid, threadsperblock](d_signatures, d_distances, num_signatures, vector_length)
d_distances.copy_to_host(distances)
print(distances[:16,:16])
# time python t4.py
[[   0.  416.  832. 1248.    0.  416.  832. 1248.    0.  416.  832. 1248.    0.  416.  832. 1248.]
 [ 416.    0.  416.  832.  416.    0.  416.  832.  416.    0.  416.  832.  416.    0.  416.  832.]
 [ 832.  416.    0.  416.  832.  416.    0.  416.  832.  416.    0.  416.  832.  416.    0.  416.]
 [1248.  832.  416.    0. 1248.  832.  416.    0. 1248.  832.  416.    0. 1248.  832.  416.    0.]
 [   0.  416.  832. 1248.    0.  416.  832. 1248.    0.  416.  832. 1248.    0.  416.  832. 1248.]
 [ 416.    0.  416.  832.  416.    0.  416.  832.  416.    0.  416.  832.  416.    0.  416.  832.]
 [ 832.  416.    0.  416.  832.  416.    0.  416.  832.  416.    0.  416.  832.  416.    0.  416.]
 [1248.  832.  416.    0. 1248.  832.  416.    0. 1248.  832.  416.    0. 1248.  832.  416.    0.]
 [   0.    0.    0.    0.    0.    0.    0.    0.    0.  416.  832. 1248.    0.  416.  832. 1248.]
 [   0.    0.    0.    0.    0.    0.    0.    0.  416.    0.  416.  832.  416.    0.  416.  832.]
 [   0.    0.    0.    0.    0.    0.    0.    0.  832.  416.    0.  416.  832.  416.    0.  416.]
 [   0.    0.    0.    0.    0.    0.    0.    0. 1248.  832.  416.    0. 1248.  832.  416.    0.]
 [   0.    0.    0.    0.    0.    0.    0.    0.    0.  416.  832. 1248.    0.  416.  832. 1248.]
 [   0.    0.    0.    0.    0.    0.    0.    0.  416.    0.  416.  832.  416.    0.  416.  832.]
 [   0.    0.    0.    0.    0.    0.    0.    0.  832.  416.    0.  416.  832.  416.    0.  416.]
 [   0.    0.    0.    0.    0.    0.    0.    0. 1248.  832.  416.    0. 1248.  832.  416.    0.]]

real    0m11.580s
user    0m5.579s
sys     0m5.922s
#

В этом случае я ограничил *Всего от 1024 * векторов до 60000 вместо 60320, поскольку большее число приводило к тому, что выходная матрица была слишком большой, чтобы поместиться в доступную память моего 16-ГБ графического процессора Tesla P100. Если в вашем графическом процессоре меньше 16 ГБ памяти, этот код не будет работать как есть. Вам нужно будет уменьшить проблему до меньшего числа векторов сигнатур.Было бы относительно просто разделить векторы на две две группы и выполнить вычисления расстояния между двумя группами, чтобы заполнить всю матрицу.

Однако код numba на моей тестовой машине выполняется примерно за 11 секунд,и, кажется, дает эквивалентный результат для очень маленькой 16x16 части выходной матрицы, которую я распечатываю.

Если я профилирую этот код, я обнаруживаю, что ядро ​​GPU работает примерно через 3 секунды,Время передачи данных огромной выходной матрицы от GPU к CPU занимает около 6 секунд, а остаток составляет около 2 секунд при загрузке Python.

Фактический алгоритм GPU ориентирован на блок.Каждый блок отвечает за сравнение 8 векторов со всем массивом векторов.Каждый блок начинается с загрузки 8 векторов в общую память, а затем пересекает весь векторный массив, вычисляя манхэттенское расстояние по каждому из этих 8 векторов.Блок использует цикл блочного шага, чтобы пройти через весь массив.В конце каждой итерации цикла в выходной массив записываются 8 манхэттенских расстояний, соответствующих текущему сравнению.

Кроме того, в коде есть небольшое изменение, так что только выходные значения матрицы вышедиагонали вычисляются.Поскольку вычисления выполняются по блокам, блоки, у которых нет элементов выше диагонали, не вычисляются.Это сокращает время обработки примерно вдвое, но полная матрица вывода все еще находится в памяти.По этой причине нижний левый квадрант моего выходного сигнала 16x16 выше равен нулю, поскольку квадрант 8x8 полностью «ниже» диагонали, поэтому обработка для него пропускается из-за ситуации сходства, уже указанной в вопросе.