---
title: "Przetwarzanie tekstu w OpenCL"
url: https://michalasobczak.pl/pl/posts/text-processing-in-opencl.html
author: "Michał Sobczak"
description: "OpenCL doskonale radzi sobie z liczbami, ale już nie tak dobrze z przetwarzaniem tekstu. Brakuje w nim nawet podstawowych funkcji dostępnych w zwykłym C99. Pytanie więc, czy warto próbować przetwarzać w nim tekst."
published: 2022-12-30T11:06:57.000Z
updated: 2024-08-29T08:23:36.000Z
language: pl
categories: ["AI/ML"]
tags: ["AI/ML", "HPC", "Programming", "OpenCL"]
license: CC BY-NC-ND 4.0 (https://creativecommons.org/licenses/by-nc-nd/4.0/)
---
# Przetwarzanie tekstu w OpenCL

OpenCL doskonale radzi sobie z liczbami, ale już nie tak dobrze z przetwarzaniem tekstu. Brakuje w nim nawet podstawowych funkcji dostępnych w zwykłym C99. Pytanie więc, czy warto próbować przetwarzać w nim tekst.

Do mojego bazowego projektu OpenCL (który można znaleźć [tutaj](https://github.com/michalasobczak/publications/tree/master/opencl/aiml/src/com/michalasobczak/aiml)) dodałem moduł „aiml”. Wczytuje on ponad 31 tys. linii tekstu zawierających ponad 4 mln znaków. Sam tekst znajduje się w pierwszym buforze, w tablicy `uchar`. Drugi bufor przechowuje wskaźniki i długości kolejnych linii, które są work-itemami, więc takich work-itemów jest ponad 31 tys. Trzeci bufor to tablica wyników, w której mogę zapisać rezultat przetwarzania przez kernel.

## Bazowy kod w Javie

Najpierw wczytujemy plik tekstowy i generujemy wskaźniki oraz długości:

```java
    public void readFile() {
        System.out.println(" - Read file");
        File file = new File("aiml/src/com/michalasobczak/aiml/bible.txt");
        try {
            BufferedReader br = new BufferedReader(new FileReader(file));
            String st = null;
            while ((st = br.readLine()) != null) {
                allstring.append(st);
                vector.add(st);
                line_sizes.add(pointer);
                pointer = pointer + st.length();
                counter++;
            }
        } catch (IOException e) {
            throw new RuntimeException(e);
        }
        System.out.println("Read no of lines: " + counter);
        System.out.println("sample line: " + vector.get(100));
        System.out.println("sample line size in chars: " + line_sizes.get(5));
        n = counter;
    }
```

Następnie kopiujemy tekst źródłowy, wskaźniki i długości do buforów:

```java
    public void generateSampleRandomData() {
        System.out.println(" - Started sampling data");
        for (int i=0; i<allstring.length(); i++) {
            srcArrayA[i] = (byte)allstring.charAt(i);
        }
        System.out.println("allstring size: " + allstring.length());
        for (int i=0; i<n; i++) {
            srcArrayB[i*2]  = line_sizes.get(i);
            srcArrayB[(i*2)+1] = vector.get(i).length();
        }
        System.out.println(" - Finished sampling data");
    }
```

Bufory wyglądają następująco:

```java
    public void createBuffers() {
        // Allocate the memory objects for the input- and output data
        this.memObjects[0] = clCreateBuffer(this.context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, (long) Sizeof.cl_uchar * this.allstring.length(), KernelConfigurationSet.srcA, null);
        this.memObjects[1] = clCreateBuffer(this.context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, (long) Sizeof.cl_int * this.n * 2, KernelConfigurationSet.srcB, null);
        this.memObjects[2] = clCreateBuffer(this.context, CL_MEM_READ_WRITE, (long) Sizeof.cl_uchar * this.n, null, null);
    }
```

Kod kernela jest podzielony na dwie części:

```java
    public void readKernelFile() {
        this.content = new String("");
        try {
            this.content = Files.readString(Path.of("aiml/src/com/michalasobczak/aiml/utils.c"));
            this.content += "\n";
            this.content += Files.readString(Path.of("aiml/src/com/michalasobczak/aiml/kernel.c"));

        } catch (IOException e) {
            e.printStackTrace();
        }
    }
```

Jeśli chodzi o bazowy kod w Javie, to wszystko.

## Kod kernela w C99

Jak już wspomniałem, w OpenCL C99 brakuje kilku podstawowych funkcji do przetwarzania tekstu. Myślę, że po pierwsze dlatego, że nie jest przeznaczony do pracy z tekstem, a po drugie dlatego, że nie można stosować rekurencji, a podejrzewam, że niektóre z tych funkcji mogą z niej korzystać. Postanowiłem więc przygotować kilka podstawowych funkcji:

```cpp
#define MAX_STRING_SIZE 256

int get_string_size(uchar string[]) {
    int size;
    int c;
    int counter = 0;
    for (c=0; c<MAX_STRING_SIZE; c++) {
        if (string[c] == NULL) {
            break;
        }
        counter++;
    }
    return counter;
}

void print_string(uchar stringa[]) {
    int size = get_string_size(stringa);
    int c;
    for (c=0; c<size; c++) {
        printf("%c", stringa[c]);
    }
    printf("\n");
}

void set_string(uchar wb[], uchar tmp[]) {
    int c;
    int size = get_string_size(tmp);
    for (c=0; c<size; c++) {
        wb[c] = tmp[c];
    }
}

void set_string_and_print(uchar wb[], uchar tmp[]) {
    set_string(wb, tmp);
    print_string(wb);
}

int find_string_in_string(uchar source[], uchar looking_for[]) {
    int s_size = get_string_size(source);
    int lf_size = get_string_size(looking_for);
    int c, d;
    for (c=0; c<s_size; c++) {
        for (d=0; d<lf_size; d++) {
            if (source[c+d] == looking_for[d]) {
                ;
            }
            else {
                break;
            }
            if (d == lf_size-1) {
                return 1;
            }
        } 
    } 
    return 0;
}
```

Kilka słów wyjaśnienia. Funkcja zwracająca rozmiar łańcucha opiera się na tablicy znaków zakończonej znakiem NULL. Funkcja ustawiająca łańcuch nie dodaje tego NULL na końcu, więc w razie potrzeby musisz zrobić to sam. Wyszukiwanie łańcucha w łańcuchu zwraca tylko pierwsze trafienie.

A teraz kernel:

```cpp
__kernel void sampleKernel(__global const uchar* srcA,
                           __global const int2*   srcB,
                           __global       uchar*  dst)
{
    int gc, ls, gs, gid, lid;
    int c, d;
    int res = 0;
    __private uchar current[131*1000]    = { '\0' };
    __private uchar word_buffer[30] = { '\0' };
    // -- PROCEDURE
    gid = get_global_id(0);
    int2 val = srcB[gid];
    d = 0;
    // prepare current local item
    for (c=val.s0; c<val.s0+val.s1; c++) {
        current[d] = srcA[c];
        d++;
    } // for
    uchar tmp[10] = "LORD\0";
    set_string(word_buffer, tmp);
    res = find_string_in_string(current, word_buffer);
    dst[gid] = res;
} // kernel
```

Jak pokazałem wcześniej w bazowym kodzie w Javie, są 3 bufory. Pierwszy to zwykła tablica `uchar` z tekstem źródłowym. Drugi zawiera wektory `int2` przechowujące wskaźniki i długości kolejnych, sąsiadujących linii/wersetów. Trzeci jest na dane wyjściowe – na przykład w przypadku udanego wyszukiwania zawiera 1, w przeciwnym razie 0.

Przetestowałem to na moim NVIDIA GeForce RTX 3050 Ti Mobile z 4 GB VRAM. Około 32 tys. elementów (work-itemów) oznacza, że na pojedynczy work-item możemy przydzielić aż 131kB danych. W ten sposób możemy w pełni wykorzystać całą dostępną pamięć VRAM. Oczywiście nie wszystkie work-itemy zostaną uruchomione jednocześnie, ponieważ to GPU ma tylko 2560 rdzeni. To więc, rzecz jasna, maksymalna liczba elementów przetwarzanych równolegle „w tym samym czasie”. Szacunkowo do przetworzenia wszystkich danych potrzeba 13 rund, musimy jednak pamiętać, że local-size jest ustawiony na 32 i że specyfikacja CC (compute capabilities) nakłada na samo GPU pewne specyficzne ograniczenia.

Dla CC 8.6 mamy maksymalnie 16 bloków wątków na SM (streaming multiprocessor), co pomnożone przez 32 work-itemy local-work-size daje maksymalnie 512. RTX 3050 Ti ma 20 SM, więc maksymalna liczba jednocześnie (teoretycznie) przetwarzanych elementów wynosiłaby 10240, ale przy zaledwie 2560 rdzeniach sądzę, że oczywiście nie dojdzie aż tak daleko i osiągnie 100% wykorzystania przy dużo niższych wartościach. Mimo to najnowsze GPU mogą mieć do 16 tys. rdzeni, więc taki sprzęt mógłby lepiej wykorzystać CC 8.6 lub wyższą specyfikację przy pełnym obciążeniu.

Chciałbym zwrócić uwagę na jeszcze jedną rzecz dotyczącą pamięci prywatnej/rejestrów i pamięci globalnej. W przypadku architektury GPU Ampere:

-   Rozmiar pliku rejestrów wynosi 64K 32-bitowych rejestrów na SM.
-   Maksymalna liczba rejestrów na wątek wynosi 255.

Jesteśmy więc ograniczeni do 255 rejestrów na work-item, a do tego obowiązuje limit 64k na SM. Możemy zatem oszacować, a nawet obliczyć maksymalny rozmiar danych, który zmieści się lokalnie – powyżej tej wartości dane trafią do pamięci globalnej, co oznacza znacznie większe opóźnienia. Widać to po rosnących czasach obliczeń, gdy zwiększamy tablicę `uchar current`.

## Podsumowanie

Przetwarzanie tekstu w OpenCL działa bez zarzutu – 0 – 1 ms na pojedyncze wyszukiwanie w ponad 4 mln znaków (31-32 tys. linii). Ogranicza nas brak funkcji do obsługi łańcuchów i pamięci, więc wszystkie funkcje łańcuchowe, które napisałem, korzystają z buforów w postaci stałych tablic. W praktyce przetestowałem przydział całej pamięci VRAM (4GB). Pobór mocy wynosi 51W.

![](https://michalasobczak.pl/wp-content/uploads/2022/12/Screenshot-from-2022-12-30-12-02-18.png)

Co dalej? Myślę, że mogę pójść krok dalej i spróbować zrobić jakąś klasyfikację oraz kilka innych rzeczy w kierunku ML, a nawet AI. To może być całkiem ciekawe…