字符串數(shù)據(jù)的高效處理對于許多數(shù)據(jù)科學應(yīng)用至關(guān)重要。為了從字符串數(shù)據(jù)中提取有價值的信息， RAPIDS libcudf 提供了加速字符串數(shù)據(jù)轉(zhuǎn)換的強大工具。 lib cuDF 是一個 C ++ GPU DataFrame 庫，用于加載、連接、聚合和過濾數(shù)據(jù)。

在數(shù)據(jù)科學中，字符串數(shù)據(jù)表示語音、文本、遺傳序列、日志記錄和許多其他類型的信息。在為機器學習和特征工程處理字符串數(shù)據(jù)時，必須經(jīng)常對數(shù)據(jù)進行規(guī)范化和轉(zhuǎn)換，然后才能將其應(yīng)用于特定用例。 lib cuDF 提供了通用 API 和設(shè)備端實用程序，以支持廣泛的自定義字符串操作。

本文演示了如何使用 lib cuDF 通用 API 巧妙地轉(zhuǎn)換字符串列。您將獲得關(guān)于如何使用自定義內(nèi)核和 lib cuDF 設(shè)備端實用程序解鎖峰值性能的新知識。本文還向您介紹了如何最好地管理 GPU 內(nèi)存和有效地構(gòu)造 lib cuDF 列以加快字符串轉(zhuǎn)換的示例。

介紹字符串列的箭頭格式

lib cuDF 使用 Arrow format 將字符串數(shù)據(jù)存儲在設(shè)備內(nèi)存中，它將字符串列表示為兩個子列： chars and offsets （圖 1 ）。

chars列將字符串數(shù)據(jù)保存為 UTF-8 編碼的字符字節(jié)，并連續(xù)存儲在內(nèi)存中。

offsets列包含遞增的整數(shù)序列，這些整數(shù)是字節(jié)位置，用于標識字符數(shù)據(jù)數(shù)組中每個單獨字符串的開頭。最后一個 offset 元素是 chars 列中的字節(jié)總數(shù)。這意味著第i行的單個字符串的大小定義為（offsets[i+1]-offsets[i])。

圖 1.顯示箭頭格式如何用chars和offsets子列表示字符串列的示意圖

字符串編校函數(shù)示例

要演示字符串轉(zhuǎn)換的示例，請考慮一個函數(shù)，該函數(shù)接收兩個輸入字符串列，并生成一個經(jīng)過修訂的輸出字符串列。

輸入數(shù)據(jù)的格式如下：一個“名稱”列包含由空格分隔的名字和姓氏，另一個“可見性”列包含“公共”或“私有”狀態(tài)

我們建議使用“ redact ”函數(shù)對輸入數(shù)據(jù)進行操作，以生成由姓氏的首字母加上空格和整個名字組成的輸出數(shù)據(jù)。但是，如果相應(yīng)的可見性列為“ private ”，則輸出字符串應(yīng)完全修訂為“ X X ”

表 1.“編校”字符串轉(zhuǎn)換的示例，該轉(zhuǎn)換接收名稱和可見性字符串列作為輸入，部分或完全編校的數(shù)據(jù)作為輸出

使用 lib cuDF API 轉(zhuǎn)換字符串

首先，可以使用 libcudf strings API 完成字符串轉(zhuǎn)換。通用 API 是比較性能的良好起點和基線。

API 函數(shù)操作整個字符串列，每個函數(shù)至少啟動一個內(nèi)核，每個字符串分配一個線程。每個線程跨 GPU 并行處理一行數(shù)據(jù)，并輸出一行作為新輸出列的一部分。

要使用通用 API 完成修訂示例函數(shù)，請執(zhí)行以下步驟：

使用contains將“可見性”字符串列轉(zhuǎn)換為布爾列

每當布爾列中相應(yīng)的行條目為“ false ”時，通過復(fù)制“ X X ”從名稱列創(chuàng)建新的字符串列

將“修訂”列拆分為名字和姓氏列

將姓氏的第一個字符切成姓氏首字母

通過用空格（“”）分隔符連接最后一個縮寫列和第一個名字列來構(gòu)建輸出列。

// convert the visibility label into a boolean
auto const visible = cudf::string_scalar(std::string("public"));
auto const allowed = cudf::strings::contains(visibilities, visible);

// redact names 
auto const redaction = cudf::string_scalar(std::string("X X"));
auto const redacted = cudf::copy_if_else(names, redaction, allowed->view());

// split the first name and last initial into two columns
auto const sv = cudf::strings_column_view(redacted->view())
auto const first_last  = cudf::strings::split(sv);
auto const first = first_last->view().column(0);
auto const last  = first_last->view().column(1);
auto const last_initial = cudf::strings::slice_strings(last, 0, 1);  

// assemble a result column
auto const tv = cudf::table_view({last_initial->view(), first});
auto result = cudf::strings::concatenate(tv, std::string(" "));

在具有 600K 行數(shù)據(jù)的 A6000 上，此方法大約需要 3.5 毫秒。本例使用contains、copy_if_else, split, slice_strings和concatenate完成自定義字符串轉(zhuǎn)換。使用 Nsight Systems 進行的分析表明，split函數(shù)花費的時間最長，其次是和concatenate。

圖 2 顯示了修訂示例的 Nsight Systems 的分析數(shù)據(jù)，顯示了每秒最多 6 億個元素的端到端字符串處理。這些區(qū)域?qū)?yīng)于與每個功能相關(guān)聯(lián)的 NVTX 范圍。淺藍色范圍對應(yīng) CUDA 內(nèi)核運行的周期。

圖 2.修訂示例中來自 Nsight Systems 的分析數(shù)據(jù)

使用自定義內(nèi)核轉(zhuǎn)換字符串

lib cuDF strings API 是一個快速高效的字符串轉(zhuǎn)換工具包，但有時性能關(guān)鍵型函數(shù)需要運行得更快。 lib cuDF strings API 中額外工作的一個關(guān)鍵來源是在全局設(shè)備內(nèi)存中為每個 API 調(diào)用創(chuàng)建至少一個新的字符串列，這為將多個 API 調(diào)用合并到自定義內(nèi)核提供了機會。

內(nèi)核 malloc 調(diào)用的性能限制

首先，我們將構(gòu)建一個自定義內(nèi)核來實現(xiàn)編校示例轉(zhuǎn)換。在設(shè)計這個內(nèi)核時，我們必須記住 lib cuDF strings 列是不可變的。

不能就地更改字符串列，因為字符字節(jié)是連續(xù)存儲的，對字符串長度的任何更改都會使偏移量數(shù)據(jù)無效。因此，redact_kernel自定義內(nèi)核通過使用 lib cuDF 列工廠來構(gòu)建offsets和chars子列，從而生成一個新的字符串列。

在第一種方法中，使用內(nèi)核內(nèi)的 malloc 調(diào)用在 dynamic device memory 中創(chuàng)建每行的輸出字符串。自定義內(nèi)核輸出是一個指向每行輸出的設(shè)備指針向量，該向量用作字符串列工廠的輸入。

自定義內(nèi)核接受 cudf::column_device_view 以訪問字符串列數(shù)據(jù)，并使用element方法返回 cudf::string_view 表示指定行索引處的字符串數(shù)據(jù)。內(nèi)核輸出是一個cudf::string_view類型的向量，它保存指向設(shè)備內(nèi)存的指針，其中包含輸出字符串和該字符串的字節(jié)大小。

cudf::string_view類類似于 std :： string _ view 類，但它是專門為 lib cuDF 實現(xiàn)的，并將固定長度的字符數(shù)據(jù)包裝在編碼為 UTF-8 的設(shè)備內(nèi)存中。它與std對應(yīng)物具有許多相同的特性（例如find和[EZX 43]函數(shù)）和限制（無空終止符）。cudf::string_view表示存儲在設(shè)備內(nèi)存中的字符序列，因此我們可以在這里使用它來記錄輸出向量的 malloc 內(nèi)存。

麥芽糖核

// note the column_device_view inputs to the kernel

__global__ void redact_kernel(cudf::column_device_view const d_names,
                              cudf::column_device_view const d_visibilities,
                              cudf::string_view redaction,
                              cudf::string_view* d_output)
{
  // get index for this thread
  auto index = threadIdx.x + blockIdx.x * blockDim.x;
  if (index >= d_names.size()) return;

  auto const visible = cudf::string_view("public", 6);

  auto const name = d_names.element(index);
  auto const vis  = d_visibilities.element(index);
  if (vis == visible) {
    auto const space_idx    = name.find(' ');
    auto const first        = name.substr(0, space_idx);
    auto const last_initial = name.substr(space_idx + 1, 1);
    auto const output_size  = first.size_bytes() + last_initial.size_bytes() + 1;
    
    char* output_ptr = static_cast(malloc(output_size));

    // build output string
    d_output[index]  = cudf::string_view{output_ptr, output_size};
    memcpy(output_ptr, last_initial.data(), last_initial.size_bytes());
    output_ptr += last_initial.size_bytes();
    *output_ptr++ = ' ';
    memcpy(output_ptr, first.data(), first.size_bytes());
  } else {
    d_output[index] = cudf::string_view{redaction.data(), redaction.size_bytes()};
  }
}

__global__ void free_kernel(cudf::string_view redaction, cudf::string_view* d_output, int count)
{
  auto index = threadIdx.x + blockIdx.x * blockDim.x;
  if (index >= count) return;

  auto ptr = const_cast(d_output[index].data());
  if (ptr != redaction.data()) free(ptr); // free everything that does match the redaction string
}

在測量內(nèi)核性能之前，這似乎是一種合理的方法。在具有 600K 行數(shù)據(jù)的 A6000 上，這種方法大約需要 108 毫秒，比上面使用 lib cuDF 字符串 API 提供的解決方案慢 30 倍以上。

redact_kernel         60.3ms
free_kernel           45.5ms
make_strings_column    0.5ms

這里的主要瓶頸是兩個內(nèi)核中的malloc/free調(diào)用。 CUDA 動態(tài)設(shè)備內(nèi)存需要同步內(nèi)核中的malloc/free調(diào)用，導(dǎo)致并行執(zhí)行退化為順序執(zhí)行。

預(yù)先分配工作內(nèi)存以消除瓶頸

在啟動內(nèi)核之前，用預(yù)先分配的工作內(nèi)存替換內(nèi)核中的malloc/free調(diào)用，從而消除malloc/free瓶頸。

對于編校示例，此示例中每個字符串的輸出大小不應(yīng)大于輸入字符串本身，因為邏輯只刪除字符。因此，可以使用與輸入緩沖區(qū)大小相同的單個設(shè)備內(nèi)存緩沖區(qū)。使用輸入偏移定位每行位置。

訪問字符串列的偏移量涉及使用cudf::strings_column_view包裝cudf::column_view并調(diào)用其 offsets_begin方法。也可以使用chars_size方法訪問chars子列的大小。然后在調(diào)用內(nèi)核以存儲字符輸出數(shù)據(jù)之前預(yù)先分配rmm::device_uvector。

auto const scv = cudf::strings_column_view(names);
auto const offsets = scv.offsets_begin();
auto working_memory = rmm::device_uvector(scv.chars_size(), stream);

預(yù)分配內(nèi)核

__global__ void redact_kernel(cudf::column_device_view const d_names,
                              cudf::column_device_view const d_visibilities,
                              cudf::string_view redaction,
                              char* working_memory,
                              cudf::offset_type const* d_offsets,
                              cudf::string_view* d_output)
{
  auto index = threadIdx.x + blockIdx.x * blockDim.x;
  if (index >= d_names.size()) return;

  auto const visible = cudf::string_view("public", 6);

  auto const name = d_names.element(index);
  auto const vis  = d_visibilities.element(index);
  if (vis == visible) {
    auto const space_idx    = name.find(' ');
    auto const first        = name.substr(0, space_idx);
    auto const last_initial = name.substr(space_idx + 1, 1);
    auto const output_size  = first.size_bytes() + last_initial.size_bytes() + 1;

    // resolve output string location
    char* output_ptr = working_memory + d_offsets[index];
    d_output[index]  = cudf::string_view{output_ptr, output_size};

    // build output string into output_ptr
    memcpy(output_ptr, last_initial.data(), last_initial.size_bytes());
    output_ptr += last_initial.size_bytes();
    *output_ptr++ = ' ';
    memcpy(output_ptr, first.data(), first.size_bytes());
  } else {
    d_output[index] = cudf::string_view{redaction.data(), redaction.size_bytes()};
  }
}

內(nèi)核輸出cudf::string_view對象的矢量，該矢量被傳遞給 cudf::make_strings_column 工廠功能。此函數(shù)的第二個參數(shù)用于標識輸出列中的空條目。本文中的示例沒有空條目，因此使用了 nullptr 占位符cudf::string_view{nullptr,0}。

auto str_ptrs = rmm::device_uvector(names.size(), stream);

redact_kernel<<>>(*d_names,
                                                         *d_visibilities,
                                                         d_redaction.value(),
                                                         working_memory.data(),
                                                         offsets,
                                                         str_ptrs.data());

auto result = cudf::make_strings_column(str_ptrs, cudf::string_view{nullptr,0}, stream);

在具有 600K 行數(shù)據(jù)的 A6000 上，此方法大約需要 1.1 毫秒，因此比基線高出 2 倍以上。大致細分如下：

  redact_kernel            66us
  make_strings_column     400us

剩余時間花在cudaMalloc, cudaFree, cudaMemcpy,上，這是管理rmm::device_uvector臨時實例的典型開銷。如果保證所有輸出字符串的大小都與輸入字符串相同或更小，則此方法很有效。

總的來說，使用 RAPIDS RMM 切換到批量工作內(nèi)存分配是一個顯著的改進，也是自定義字符串函數(shù)的一個很好的解決方案。

優(yōu)化列創(chuàng)建以加快計算時間

有沒有辦法進一步改進這一點？現(xiàn)在的瓶頸是cudf::make_strings_column工廠函數(shù)，它從cudf::string_view對象的矢量構(gòu)建兩個字符串列組件offsets和chars。

在 libcudf 中，包含了許多工廠函數(shù)來構(gòu)建字符串列。前面示例中使用的工廠函數(shù)獲取cudf::string_view對象的cudf::device_span，然后通過對底層字符數(shù)據(jù)執(zhí)行g(shù)ather來構(gòu)造列，以構(gòu)建偏移量和字符子列。rmm::device_uvector可自動轉(zhuǎn)換為cudf::device_span，而無需復(fù)制任何數(shù)據(jù)。

但是，如果直接構(gòu)建字符矢量和偏移矢量，則可以使用不同的工廠函數(shù)，它只需創(chuàng)建字符串列，而不需要聚集來復(fù)制數(shù)據(jù)。

sizes_kernel首先傳遞輸入數(shù)據(jù)，以計算每個輸出行的確切輸出大小：

Optimized kernel: Part 1

__global__ void sizes_kernel(cudf::column_device_view const d_names,
                             cudf::column_device_view const d_visibilities,
                             cudf::size_type* d_sizes)
{
  auto index = threadIdx.x + blockIdx.x * blockDim.x;
  if (index >= d_names.size()) return;

  auto const visible = cudf::string_view("public", 6);
  auto const redaction = cudf::string_view("X X", 3);

  auto const name = d_names.element(index);
  auto const vis  = d_visibilities.element(index);

  cudf::size_type result = redaction.size_bytes(); // init to redaction size
  if (vis == visible) {
    auto const space_idx    = name.find(' ');
    auto const first        = name.substr(0, space_idx);
    auto const last_initial = name.substr(space_idx + 1, 1);

    result = first.size_bytes() + last_initial.size_bytes() + 1;
  }

  d_sizes[index] = result;
}

然后，通過執(zhí)行就地exclusive_scan將輸出大小轉(zhuǎn)換為偏移量。請注意，offsets矢量是用names.size()+1元素創(chuàng)建的。最后一項是字節(jié)總數(shù)（所有大小加在一起），而第一項是 0 。這兩項都由exclusive_scan調(diào)用處理。從offsets列的最后一個條目檢索chars列的大小，以構(gòu)建字符矢量。

// create offsets vector
auto offsets = rmm::device_uvector(names.size() + 1, stream);

// compute output sizes
sizes_kernel<<>>(
  *d_names, *d_visibilities, offsets.data());

thrust::exclusive_scan(rmm::exec_policy(stream), offsets.begin(), offsets.end(), offsets.begin());

redact_kernel邏輯仍然非常相同，只是它接受輸出d_offsets矢量來解析每行的輸出位置：

優(yōu)化內(nèi)核：第 2 部分

__global__ void redact_kernel(cudf::column_device_view const d_names,
                              cudf::column_device_view const d_visibilities,
                              cudf::size_type const* d_offsets,
                              char* d_chars)
{
  auto index = threadIdx.x + blockIdx.x * blockDim.x;
  if (index >= d_names.size()) return;

  auto const visible = cudf::string_view("public", 6);
  auto const redaction = cudf::string_view("X X", 3);

  // resolve output_ptr using the offsets vector
  char* output_ptr   = d_chars + d_offsets[index];

  auto const name = d_names.element(index);
  auto const vis = d_visibilities.element(index);
  if (vis == visible) {
    auto const space_idx    = name.find(' ');
    auto const first        = name.substr(0, space_idx);
    auto const last_initial = name.substr(space_idx + 1, 1);
    auto const output_size  = first.size_bytes() + last_initial.size_bytes() + 1;

    // build output string
    memcpy(output_ptr, last_initial.data(), last_initial.size_bytes());
    output_ptr += last_initial.size_bytes();
    *output_ptr++ = ' ';
    memcpy(output_ptr, first.data(), first.size_bytes());
  } else {
    memcpy(output_ptr, redaction.data(), redaction.size_bytes());
  }
}

從d_offsets列的最后一個條目檢索輸出d_chars列的大小，以分配字符矢量。內(nèi)核使用預(yù)先計算的偏移量向量啟動，并返回填充的字符向量。最后， lib cuDF strings 列工廠創(chuàng)建輸出字符串列。

這 cudf::make_strings_column factory 函數(shù)構(gòu)建 strings 列而不復(fù)制數(shù)據(jù)。offsets數(shù)據(jù)和 chars數(shù)據(jù)已經(jīng)采用了正確的預(yù)期格式，該工廠只需從每個矢量中移動數(shù)據(jù)并圍繞它創(chuàng)建列結(jié)構(gòu)。完成后，offsets和chars的rmm::device_uvectors為空，它們的數(shù)據(jù)已移動到輸出列中。

cudf::size_type output_size = offsets.back_element(stream);
auto chars = rmm::device_uvector(output_size, stream);

redact_kernel<<>>(
    *d_names, *d_visibilities, offsets.data(), chars.data());


// from pre-assembled offsets and character buffers
auto result = cudf::make_strings_column(names.size(), std::move(offsets), std::move(chars));

在具有 600K 行數(shù)據(jù)的 A6000 上，此方法需要大約 300 us （ 0.3 ms ）的時間，比以前的方法提高了 2 倍多。您可能會注意到sizes_kernel和redact_kernel共享很多相同的邏輯：一次測量輸出的大小，然后再次填充輸出。

從代碼質(zhì)量的角度來看，將轉(zhuǎn)換重構(gòu)為由大小和編校內(nèi)核調(diào)用的設(shè)備函數(shù)是有益的。從性能的角度來看，您可能會驚訝地看到轉(zhuǎn)換的計算成本是原來的兩倍。

內(nèi)存管理和更高效的列創(chuàng)建的好處往往超過兩次執(zhí)行轉(zhuǎn)換的計算成本。

表 2 顯示了本文討論的四種解決方案的計算時間、內(nèi)核計數(shù)和處理的字節(jié)數(shù)。“內(nèi)核啟動總數(shù)”反映了啟動的內(nèi)核總數(shù)，包括計算內(nèi)核和輔助內(nèi)核。“處理的總字節(jié)數(shù)”是累積的 DRAM 讀寫吞吐量，“處理的最小字節(jié)數(shù)”為測試輸入和輸出的平均每行 37.9 字節(jié)。理想的“內(nèi)存帶寬受限”情況假設(shè) 768 GB / s 帶寬，即 A6000 的理論峰值吞吐量。

表 2.本文討論的四種解決方案的計算時間、內(nèi)核計數(shù)和處理字節(jié)

由于內(nèi)核啟動次數(shù)減少，處理的總字節(jié)數(shù)減少，“優(yōu)化內(nèi)核”提供了最高的吞吐量。有了高效的自定義內(nèi)核，內(nèi)核的總啟動次數(shù)從 31 次減少到 4 次，處理的總字節(jié)數(shù)從輸入加輸出大小的 12.6 倍減少到 1.75 倍。

因此，定制內(nèi)核的吞吐量比用于編校轉(zhuǎn)換的通用字符串 API 高出 10 倍以上。

峰值性能分析

RAPIDS Memory Manager (RMM) 中的池內(nèi)存資源是另一個可用于提高性能的工具。上述示例使用默認的“ CUDA 內(nèi)存資源”分配和釋放全局設(shè)備內(nèi)存。然而，分配工作內(nèi)存所需的時間增加了字符串轉(zhuǎn)換步驟之間的延遲。 RMM 中的“池內(nèi)存資源”通過預(yù)先分配一個大的內(nèi)存池，并在處理過程中根據(jù)需要分配子分配來減少延遲。

對于 CUDA 內(nèi)存資源，“ Optimized Kernel ”顯示了 10×15 倍的加速，由于分配大小的增加，在較高的行數(shù)時開始下降（圖 3 ）。使用池內(nèi)存資源可以緩解這種影響，并比 lib cuDF stringsAPI 方法保持 15x-25 倍的速度提升。

圖 3.使用默認 CUDA 內(nèi)存資源（實線）和池內(nèi)存資源（虛線）從自定義內(nèi)核“預(yù)分配內(nèi)核”和“優(yōu)化內(nèi)核”加速，而使用默認 CUDA 存儲資源的 lib cuDF 字符串 API

利用池內(nèi)存資源，兩遍算法的端到端內(nèi)存吞吐量接近理論極限。“優(yōu)化內(nèi)核”的吞吐量達到 320-340 GB / s ，使用輸入大小加上輸出大小和計算時間進行測量（圖 4 ）。

兩遍法首先測量輸出元素的大小，分配內(nèi)存，然后用輸出設(shè)置內(nèi)存。給定兩遍處理算法，“優(yōu)化內(nèi)核”中的實現(xiàn)性能接近內(nèi)存帶寬限制。“端到端內(nèi)存吞吐量”定義為輸入加輸出大小（ GB ）除以計算時間* RTX A6000 內(nèi)存帶寬（ 768 GB / s ）。

關(guān)鍵要點

本文演示了在 libcudf 中編寫高效字符串數(shù)據(jù)轉(zhuǎn)換的兩種方法。 lib cuDF 通用 API 對于開發(fā)人員來說是快速而直接的，并且提供了良好的性能。 lib cuDF 還提供了設(shè)計用于自定義內(nèi)核的設(shè)備端實用程序，在本例中，解鎖性能提高了 10 倍以上。