Статьи

ГПГПУ с Джудой Хорошей, Плохой и… Уродливой

Январь 5, 2019
В нашей предыдущей статье GPGPU для Java-программирования мы показали, как настроить среду для выполнения CUDA из кода Java. Однако предыдущая статья была сосредоточена только на настройке среды, оставляя предмет параллелизма нетронутым. В этой статье мы увидим, как мы можем использовать графический процессор для достижения наилучших результатов: параллельной обработки. В этом примере мы возьмем некоторые метрики и посмотрим, где обработка в GPU сильнее или слабее, чем при использовании CPU… и, конечно, как видно из заголовка, в конце есть неприятная часть.

Мы начнем наше исследование параллелизма GPU с разработки примера, который будет отличаться от примеров, найденных в большей части доступной документации GPGPU, которая в основном написана людьми с сильным графическим или научным опытом. Большинство из этих примеров говорят о векторных добавлениях или какой-либо другой математической конструкции. Давайте поработаем над примером, который несколько напоминает деловые ситуации. Итак, давайте начнем с воображения, что у нас есть список продуктов, каждый из которых имеет кодовый тег и цену, и мы хотели бы применить накладные 10% ко всем продуктам, код которых «abc».

Сначала мы реализуем пример на C, чтобы сделать некоторые измерения производительности между процессором и обработкой графического процессора. Впоследствии, конечно, мы будем реализовывать то же самое в Java, но мы будем избегать каких-либо измерений, поскольку они немного сложнее в Java, так как мы должны учитывать такие вещи, как сборка мусора, просто во время компиляции и т. Д.

001
002
003
004
005
006
007
008
009
010
011
012
013
014
015
016
017
018
019
020
021
022
023
024
025
026
027
028
029
030
031
032
033
034
035
036
037
038
039
040
041
042
043
044
045
046
047
048
049
050
051
052
053
054
055
056
057
058
059
060
061
062
063
064
065
066
067
068
069
070
071
072
073
074
075
076
077
078
079
080
081
082
083
084
085
086
087
088
089
090
091
092
093
094
095
096
097
098
099
100
101
102
103
//============================================================================
// Name        : StoreDiscountExample.cu
// Author      : Spyros Sakellariou
// Version     : 1.0
// Description : The Good the Bad and the Ugly
//============================================================================
 
#include <iostream>
#include <sys/time.h>
  
typedef struct {
  char code[3];
 float listPrice;
} product;
 
void printProducts(long size, product * myProduct){
printf("Price of First item=%f,%s\n",myProduct[0].listPrice,myProduct[0].code);
printf("Price of Second item=%f,%s\n",myProduct[1].listPrice,myProduct[1].code);
printf("Price of Middle item=%f,%s\n",myProduct[(size-1)/2].listPrice,myProduct[(size-1)/2].code);
printf("Price of Almost Last item=%f,%s\n",myProduct[size-2].listPrice,myProduct[size-2].code);
printf("Price of Last item=%f,%s\n",myProduct[size-1].listPrice,myProduct[size-1].code);
}
 
float calculateMiliseconds (timeval t1,timeval t2) {
        float elapsedTime;
 elapsedTime = (t2.tv_sec - t1.tv_sec) * 1000.0;
  elapsedTime += (t2.tv_usec - t1.tv_usec) / 1000.0;
  return elapsedTime;
}
  
  
 
__global__ void kernel(long size, product *products)
{
  long kernelid = threadIdx.x + blockIdx.x * blockDim.x;
 while(kernelid < size) {
    if (products[kernelid].code[0]=='a' && products[kernelid].code[1]=='b' && products[kernelid].code[2]=='c')
  products[kernelid].listPrice*=1.10;
    kernelid += blockDim.x * gridDim.x;
  }
}
 
int main( int argc, char** argv)
{
  timeval t1,t2;
  cudaEvent_t eStart,eStop;
  float elapsedTime;
  long threads = 256;
  long blocks = 1024;
  long size = 9000000;
  char *product1 = "abc";
  char *product2 = "bcd";
  product *myProduct;
  product *dev_Product;
 
 printf("blocks=%d x threads=%d total threads=%d total number of products=%d\n\n",blocks,threads,threads*blocks,size);
 
 myProduct=(product*)malloc(sizeof(myProduct)*size);
  cudaMalloc((void**)&dev_Product,sizeof(dev_Product)*size);
 cudaEventCreate(&eStart);
 cudaEventCreate(&eStop);
  
 gettimeofday(&t1, NULL);
  for (long i = 0; i<size; i++){
   if (i%2==0)
    strcpy(myProduct[i].code,product1);
   else
   strcpy(myProduct[i].code,product2);
  myProduct[i].listPrice = i+1;
 }
  gettimeofday(&t2, NULL);
  printf ( "Initialization time %4.2f ms\n", calculateMiliseconds(t1,t2) );
 printProducts(size,myProduct);
  cudaMemcpy(dev_Product,myProduct,sizeof(dev_Product)*size,cudaMemcpyHostToDevice);
   
 cudaEventRecord(eStart,0);
  kernel<<<blocks,threads>>>(size,dev_Product);
  cudaEventRecord(eStop,0);
  cudaEventSynchronize(eStop);
   
  cudaMemcpy(myProduct,dev_Product,sizeof(dev_Product)*size,cudaMemcpyDeviceToHost);
   
  cudaEventElapsedTime(&elapsedTime,eStart,eStop);
  printf ( "\nCuda Kernel Time=%4.2f ms\n", elapsedTime );
  printProducts(size,myProduct);
  
  
  gettimeofday(&t1, NULL);
  long j=0;
   while (j < size){
     if (myProduct[j].code[0]=='a' && myProduct[j].code[1]=='b' && myProduct[j].code[2]=='c')
        myProduct[j].listPrice*=0.5;
     j++;
  
  }
  gettimeofday(&t2, NULL);
  
  printf ( "\nCPU Time=%4.2f ms\n", calculateMiliseconds(t1,t2) );
  printProducts(size,myProduct);
  
  cudaFree(dev_Product);
  free(myProduct);
}

В строках 11-14 приведено определение структуры, содержащей наш продукт, с массивом символов для кода продукта и с плавающей точкой для его цены.

В строках 16 и 24 есть определение двух служебных методов, один из которых печатает некоторые продукты (поэтому мы видим, была ли выполнена работа), а другой — для преобразования необработанных разностей дат в миллисекунды. Обратите внимание, что использование стандартной функции C clock не будет работать, поскольку ее детализация недостаточна для измерения миллисекунд.

В строке 33 написано наше ядро. По сравнению с предыдущей статьей это выглядит несколько сложнее, поэтому давайте разберем его дальше …

В строке 35 мы определяем параметр kernelid . Этот параметр будет содержать уникальный идентификатор потока выполняемого потока. CUDA назначает каждому потоку идентификатор потока и номер идентификатора блока, который уникален только для его собственного измерения. В нашем примере мы инструктируем графический процессор для запуска 256 потоков и 1024 блоков, поэтому фактически графический процессор будет выполнять 262144 потоков. Это хорошо! Хотя CUDA предоставляет нам параметры threadIdx.x и blockIdx.x во время выполнения, нам нужно вручную создать уникальный идентификатор потока, чтобы узнать, в каком потоке мы находимся в данный момент. Уникальный идентификатор потока должен начинаться с 0 до 262143, таким образом, мы можем легко создать его, умножив число потоков на блок, который должен быть выполнен (используя параметр CUDA blockDim.x ) с текущим блоком, добавив его в текущий поток, таким образом:

уникальный идентификатор потока = текущий идентификатор потока + идентификатор текущего блока * количество потоков в блоке

Если вам нравится читать дальше, вы уже поняли, что, хотя 262 тысячи потоков впечатляют, наш набор данных состоит из 9 миллионов элементов, поэтому наши потоки должны обрабатывать более одного элемента одновременно. Мы делаем это, устанавливая цикл в строке 36, который проверяет, не превышает ли наш идентификатор потока наш массив данных продуктов. Цикл использует идентификатор потока в качестве индекса, но мы увеличиваем его по следующей формуле:

приращение индекса + = потоков на блок * общее количество блоков

Таким образом, каждый поток будет выполнять каждый цикл 9 миллионов / 262 тысячи раз, то есть он будет обрабатывать около 34 элементов.

Остальная часть кода ядра довольно проста и понятна: всякий раз, когда мы находим продукт с кодом «abc», мы умножаем его на 1,1 (наши 10% накладные расходы). Обратите внимание, что функция strcpy не может использоваться внутри нашего ядра. Вы получите ошибку компиляции, если попробуете. Не так хорошо, как хотелось бы!

Переходя к основной функции в строках 45 и 46, мы определяем два таймера C ( t1 и t2 ) и два таймера событий CUDA ( eStart и eStop ). Нам нужны таймеры CUDA, потому что ядро ​​выполняется асинхронно, и наша функция ядра мгновенно возвращается, и наши таймеры будут измерять только время, необходимое для завершения вызова функции. Тот факт, что код ядра возвращается мгновенно, означает, что мы позволяем процессору выполнять другие задачи во время выполнения кода GPU. Это тоже хорошо!

Параметры, следующие за нашими таймерами, говорят сами за себя: мы определяем количество потоков, блоков, размер массива продуктов и т. Д. Указатель myproduct будет использоваться для обработки ЦП, а указатель dev_product для обработки GPU.

В строках с 58 по 61 мы выделяем оперативную и графическую память для myproduct и dev_product, а также создаем таймеры CUDA, которые помогут нам измерить время выполнения ядра.

В строках с 63 по 73 мы инициализируем myproduct с кодами и ценами и печатаем время, которое потребовалось процессору для выполнения задачи. Мы также распечатываем некоторые образцы продукции из нашего массива, чтобы убедиться, что работа сделана правильно.

В строках с 74 по 85 мы копируем массив products в память GPU, запускаем наше ядро ​​с указанием количества потоков и блоков, которые мы хотим выполнить, и копируем результаты обратно в массив myproduct . Мы печатаем время, необходимое для запуска ядра и некоторых примеров продуктов, чтобы убедиться, что мы снова выполнили свою работу правильно.

Наконец, в строках с 88 по 99 мы позволяем процессору выполнять процесс, аналогичный тому, что выполнял графический процессор, то есть применять скидку 50% на все продукты, к которым добавлены накладные расходы графического процессора. Мы печатаем время, необходимое для выполнения задачи ЦП, и печатаем некоторые образцы продуктов, чтобы убедиться, что работа выполнена.

Давайте скомпилируем и запустим этот код:

01
02
03
04
05
06
07
08
09
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
# nvcc StoreDiscountExample.cu -o StoreDiscountExample
# ./StoreDiscountExample
blocks=1024 x threads=256 total threads=262144 total number of products=9000000
 
Initialization time 105.81 ms
Price of First item=1.000000,abc
Price of Second item=2.000000,bcd
Price of Middle item=4500000.000000,bcd
Price of Almost Last item=8999999.000000,abc
Price of Last item=9000000.000000,bcd
 
Cuda Kernel Time=1.38 ms
Price of First item=1.100000,abc
Price of Second item=2.000000,bcd
Price of Middle item=4500000.000000,bcd
Price of Almost Last item=9899999.000000,abc
Price of Last item=9000000.000000,bcd
 
CPU Time=59.58 ms
Price of First item=0.550000,abc
Price of Second item=2.000000,bcd
Price of Middle item=4500000.000000,bcd
Price of Almost Last item=4949999.500000,abc
Price of Last item=9000000.000000,bcd
#

Вот Это Да! Графическому процессору потребовалось 1,38 миллисекунды, чтобы сделать то, что ЦП заняло 59,58 миллисекунды (числа будут варьироваться в зависимости от вашего оборудования, конечно). Это хорошо!

Задержи лошадей на секунду! Прежде чем вы решите удалить весь свой код и начать переписывать все в CUDA, есть одна загвоздка: мы пропустили что-то серьезное, и это измеряет, сколько времени занимает копирование 9 миллионов записей из ОЗУ в память GPU и обратно. Вот код из строк с 74 по 85, измененный, чтобы иметь таймеры для измерения копирования списка продуктов в память GPU и из нее:

01
02
03
04
05
06
07
08
09
10
11
12
13
14
gettimeofday(&t1, NULL);
cudaMemcpy(dev_Product,myProduct,sizeof(dev_Product)*size,cudaMemcpyHostToDevice);
 
cudaEventRecord(eStart,0);
kernel<<<blocks,threads>>>(size,dev_Product);
cudaEventRecord(eStop,0);
cudaEventSynchronize(eStop);
 
cudaMemcpy(myProduct,dev_Product,sizeof(dev_Product)*size,cudaMemcpyDeviceToHost);
gettimeofday(&t2, NULL);
printf ( "\nCuda Total Time=%4.2f ms\n", calculateMiliseconds(t1,t2));
cudaEventElapsedTime(&elapsedTime,eStart,eStop);
printf ( "Cuda Kernel Time=%4.2f ms\n", elapsedTime );
printProducts(size,myProduct);

Давайте скомпилируем и запустим этот код:

01
02
03
04
05
06
07
08
09
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
# nvcc StoreDiscountExample.cu -o StoreDiscountExample
# ./StoreDiscountExample
blocks=1024 x threads=256 total threads=262144 total number of products=9000000
 
Initialization time 108.31 ms
Price of First item=1.000000,abc
Price of Second item=2.000000,bcd
Price of Middle item=4500000.000000,bcd
Price of Almost Last item=8999999.000000,abc
Price of Last item=9000000.000000,bcd
 
Cuda Total Time=55.13 ms
Cuda Kernel Time=1.38 ms
Price of First item=1.100000,abc
Price of Second item=2.000000,bcd
Price of Middle item=4500000.000000,bcd
Price of Almost Last item=9899999.000000,abc
Price of Last item=9000000.000000,bcd
 
CPU Time=59.03 ms
Price of First item=0.550000,abc
Price of Second item=2.000000,bcd
Price of Middle item=4500000.000000,bcd
Price of Almost Last item=4949999.500000,abc
Price of Last item=9000000.000000,bcd
 
#

Обратите внимание, что общее время CUDA составляет 55 миллисекунд, это всего на 4 миллисекунды быстрее, чем при использовании процессора в одном потоке. Это плохо!

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

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

Вот код для части Java:

01
02
03
04
05
06
07
08
09
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
import static jcuda.driver.JCudaDriver.*;
import jcuda.*;
import jcuda.driver.*;
import jcuda.runtime.JCuda;
 
 
public class StoreDiscountExample {
 public static void main(String[] args) {
  int threads = 256;
 int blocks = 1024;
 final int size = 9000000;
  byte product1[] = "abc".getBytes();
 byte product2[] = "bcd".getBytes();
  byte productList[] = new byte[size*3];
  float productPrices[] = new float[size];
  long size_array[] = {size};
    
  cuInit(0);
  CUcontext pctx = new CUcontext();
  CUdevice dev = new CUdevice();
  cuDeviceGet(dev, 0);
  cuCtxCreate(pctx, 0, dev);
  CUmodule module = new CUmodule();
  cuModuleLoad(module, "StoreDiscountKernel.ptx");
  CUfunction function = new CUfunction();
  cuModuleGetFunction(function, module, "kernel");
  
   
  int j=0;
  for (int i = 0; i<size; i++){
   j=i*3;
  if (i%2==0) {
    productList[j]=product1[0];
    productList[j+1]=product1[1];
    productList[j+2]=product1[2];
   }
   else {
    productList[j]=product2[0];
    productList[j+1]=product2[1];
    productList[j+2]=product2[2];
   }
      
    productPrices[i] = i+1;
     
   }
    
  printSamples(size, productList, productPrices);
    
  CUdeviceptr size_dev = new CUdeviceptr();
  cuMemAlloc(size_dev, Sizeof.LONG);
  cuMemcpyHtoD(size_dev, Pointer.to(size_array), Sizeof.LONG);
    
  CUdeviceptr productList_dev = new CUdeviceptr();
  cuMemAlloc(productList_dev, Sizeof.BYTE*3*size);
  cuMemcpyHtoD(productList_dev, Pointer.to(productList), Sizeof.BYTE*3*size);
    
  CUdeviceptr productPrice_dev = new CUdeviceptr();
  cuMemAlloc(productPrice_dev, Sizeof.FLOAT*size);
  cuMemcpyHtoD(productPrice_dev, Pointer.to(productPrices), Sizeof.FLOAT*size); 
    
  Pointer kernelParameters = Pointer.to(
   Pointer.to(size_dev),
   Pointer.to(productList_dev),
   Pointer.to(productPrice_dev)
  );
    
  cuLaunchKernel(function,
   blocks, 1, 1,
   threads, 1, 1,
   0, null,
   kernelParameters, null);
    
  cuMemcpyDtoH(Pointer.to(productPrices), productPrice_dev, Sizeof.FLOAT*size);
    
  printSamples(size, productList, productPrices);
   
  JCuda.cudaFree(productList_dev);
  JCuda.cudaFree(productPrice_dev);
  JCuda.cudaFree(size_dev);
 }
   
 
 public static void printSamples(int size, byte[] productList, float[] productPrices) {
   System.out.print(String.copyValueOf(new String(productList).toCharArray(), 0, 3));System.out.println(" "+productPrices[0]);
  System.out.print(String.copyValueOf(new String(productList).toCharArray(), 3, 3));System.out.println(" "+productPrices[1]);
  System.out.print(String.copyValueOf(new String(productList).toCharArray(), 6, 3));System.out.println(" "+productPrices[2]);
  System.out.print(String.copyValueOf(new String(productList).toCharArray(), 9, 3));System.out.println(" "+productPrices[3]);
  System.out.print(String.copyValueOf(new String(productList).toCharArray(), (size-2)*3, 3));System.out.println(" "+productPrices[size-2]);
  System.out.print(String.copyValueOf(new String(productList).toCharArray(), (size-1)*3, 3));System.out.println(" "+productPrices[size-1]);
 }
}

Начиная со строк 14, 15 и 16, мы видим, что мы больше не можем использовать структуру или класс продукта. На самом деле в jcuda все, что будет передаваться в качестве параметра в ядро, должно быть одномерным массивом. Таким образом, в строке 14 мы создаем двумерный массив байтов, представленный в виде одномерного массива. Размер массива списка продуктов равен числу продуктов, умноженному на размер в байтах каждого кода продукта (в нашем случае это всего три байта). Мы также создаем второй массив для хранения цен товаров в формате с плавающей запятой, и, наконец, размер нашего списка товаров также необходимо поместить в одномерный массив.
Я думаю, что к настоящему времени вы, наверное, уже догадались, что я собираюсь сказать: это просто ужасно!

В строках 29–45 мы заполняем наш список продуктов и массивы цен на продукты, а затем передаем их ядру для обработки, создавая указатели устройств CUDA, выделяя память и копируя данные в память графического процессора перед вызовом функции ядра.

Поскольку нам пришлось преобразовать все в одномерные массивы примитивов, наш код ядра также должен немного измениться:

01
02
03
04
05
06
07
08
09
10
11
12
13
extern "C"
 
__global__ void kernel(long *size, char *productCodes, float *productPrices)
{
 long kernelid = threadIdx.x + blockIdx.x * blockDim.x;
 long charIndex = kernelid*3;
 while(kernelid < size[0]) {
    if (productCodes[charIndex]=='a' && productCodes[charIndex+1]=='b' && productCodes[charIndex+2]=='c')
      productPrices[kernelid]*=1.10;
         kernelid += blockDim.x * gridDim.x;
         charIndex = kernelid*3;
        }
}

Единственное отличие состоит в том, что мы умножаем индекс kernelid на 3, чтобы найти правильный начальный символ в нашем массиве productCodes.

Давайте скомпилируем и запустим пример Java:

01
02
03
04
05
06
07
08
09
10
11
12
13
14
15
16
# nvcc -ptx StoreDiscountKernel.cu -o StoreDiscountKernel.ptx
# javac -cp ~/GPGPU/jcuda/JCuda-All-0.4.0-beta1-bin-linux-x86_64/jcuda-0.4.0-beta1.jar StoreDiscountExample.java
#java -cp ~/GPGPU/jcuda/JCuda-All-0.4.0-beta1-bin-linux-x86_64/jcuda-0.4.0-beta1.jar StoreDiscountExample
abc 1.0                                                                                                                                                                               
bcd 2.0                                                                                                                                                                               
abc 3.0                                                                                                                                                                               
bcd 4.0                                                                                                                                                                               
abc 8999999.0
bcd 9000000.0
abc 1.1
bcd 2.0
abc 3.3
bcd 4.0
abc 9899999.0
bcd 9000000.0
#

Несмотря на то, что код ужасен, он работает так же, как в C.

Итак, вот краткий обзор нашего опыта работы с GPU и jcuda:

ХОРОШО : Очень быстрая производительность (моим H / W был четырехъядерный процессор AMD Phenom II 3,4 ГГц и NVIDIA Geforce GTX 560 с 336 ядрами)

ХОРОШО : асинхронная операция, позволяющая процессору выполнять другие задачи

ПЛОХО : копии памяти накладывают существенное снижение производительности

UGLY : Jcuda, несомненно, полезен, если вы хотите запускать ядра CUDA из Java, но преобразовывать все как одномерные массивы примитивов действительно не удобно.

В нашей предыдущей статье было несколько очень интересных комментариев об инструментах Java для OpenCL (альтернатива CUDA для GPGPU). В следующей статье мы рассмотрим эти инструменты и посмотрим, выглядят ли они «красивее», чем Jcuda.

Справка: GPGPU с Jcuda the Good, the Bad и … the Ugly от нашего партнера по W4G Spyros Sakellariou .

Статьи по Теме :

39
Share: