[cuDF]C++ libcudfのDeveloper Guideを全文和訳してみた。(1/2)

Last updated at 2024-08-12Posted at 2024-08-11

C++ libcudf 全文和訳

この文書は、libcudf C++コードの貢献者のためのガイドです。libcudfのベストプラクティスを文書化するために、開発者は以下の追加ファイルも参照してください。

Documentation Guide: libcudfコードの文書化に関するガイドライン。
Testing Guide: 単体テストの作成に関するガイドライン。
Benchmarking Guide: ベンチマークの作成に関するガイドライン。

概要

libcudfは、カラム指向の表形式データを処理するための、GPU加速のデータ並列アルゴリズムを提供するC++ライブラリです。libcudfは、スライス、フィルタリング、ソート、さまざまな種類の集計、およびグループ化や結合などのデータベース型操作を含むアルゴリズムを提供します。libcudfは、PythonやJavaなどの複数の言語インターフェースを介して多くのクライアントにサービスを提供します。ユーザーは、C++コードから直接libcudfを使用することもできます。

用語集

このセクションでは、libcudfで使用される用語を定義します。

Column: 単一の型のデータの配列です。テーブルと共に、libcudfで使用される基本的なデータ構造です。ほとんどのlibcudfアルゴリズムはカラムに対して操作を行います。カラムには、各要素が有効か無効か（無効）を表す有効マスクが付いている場合があります。ネストされた型のカラムがサポートされており、カラムに子カラムがある場合があります。カラムは、cuDF PythonのSeriesに相当するC++のオブジェクトです。
Element: カラム内の個々のデータ項目です。行とも呼ばれます。
Scalar: データ型の単一要素を表す型です。
Table: 同じ数の要素を持つカラムの集合です。テーブルは、cuDF PythonのDataFrameに相当します。
View: 別のオブジェクトが所有するデータへのゼロコピーアクセス（場合によってはスライスやオフセットを伴う）を提供する非所有オブジェクトです。例としては、カラムビューやテーブルビューがあります。

ディレクトリ構造とファイル名

外部/公開libcudf APIは、機能に基づいて適切にタイトル付けされたヘッダーファイルにグループ化され、cudf/cpp/include/cudf/に配置されます。たとえば、cudf/cpp/include/cudf/copying.hppには、あるカラムから別のカラムにコピーする関数に関連するAPIが含まれています。C++ヘッダーファイルを示すために、.hppファイル拡張子が使用されていることに注意してください。

外部/公開libcudf C++ APIヘッダーファイルには、内部のすべてのシンボルにCUDF_EXPORTをマークする必要があります。これは、以下に示すようにcudf名前空間にマクロを配置することで行います。名前空間のマークアップはネストできないため、cudf名前空間は単独で保持する必要があります。

cppコードをコピーする
#pragma once

namespace CUDF_EXPORT cudf {
namespace lists {

...

} // namespace lists
} // namespace CUDF_EXPORT cudf

外部APIヘッダーの命名は、APIを実装するソースファイルが含まれているフォルダーの名前と一致する必要があります。たとえば、cudf/cpp/include/cudf/copying.hppに見られるAPIの実装は、cudf/src/copyingに配置されています。同様に、APIの単体テストはcudf/tests/copying/に配置されています。

複数の翻訳単位でlibcudf内で使用される詳細な名前空間定義を含む内部APIヘッダーは、include/cudf/detailに配置する必要があります。公開C++ APIヘッダーと同様に、内部C++ APIヘッダーには、関数をテストできるようにするために、cudf名前空間にCUDF_EXPORTマークアップが必要です。

libcudf内のすべてのヘッダーには、インクルードガードとして#pragma onceを使用する必要があります。

ファイル拡張子

.hpp : C++ヘッダーファイル
.cpp : C++ソースファイル
.cu : CUDA C++ソースファイル
.cuh : CUDAデバイスコードを含むヘッダーファイル

.cuと.cuhは必要な場合にのみ使用してください。良い指標は、__device__などのnvccによってのみ認識されるシンボルの含有です。もう一つの指標は、デバイス実行ポリシーを持つThrustアルゴリズムAPIです（libcudfでは常にrmm::exec_policyを使用します）。

コードとドキュメントのスタイルおよびフォーマット

libcudfコードでは、いくつかのケースを除いて、すべての名前にスネークケースを使用します。テンプレートパラメーター、単体テスト、およびテストケース名はパスカルケース（別名アッパーキャメルケース）を使用する場合があります。ハンガリアン記法は使用しませんが、デバイスデータ変数および対応するホストコピーの名前付けでは、時折使用することがあります。プライベートメンバー変数は通常、アンダースコアで始まります。

cppコードをコピーする
template <typename IteratorType>
void algorithm_function(int x, rmm::cuda_stream_view s, rmm::device_async_resource_ref mr)
{
  ...
}

class utility_class
{
  ...
private:
  int _rating{};
  std::unique_ptr<cudf::column> _column{};
}

TYPED_TEST_SUITE(RepeatTypedTestFixture, cudf::test::FixedWidthTypes);

TYPED_TEST(RepeatTypedTestFixture, RepeatScalarCount)
{
  ...
}

C++フォーマットはclang-formatを使用して強制されます。clang-formatをマシンに設定し、cudf/cpp/.clang-format構成ファイルを使用してすべての変更されたコードに対してclang-formatを実行し、コミットする前にコードをフォーマットする必要があります。最も簡単な方法は、エディターを設定して「保存時にフォーマット」することです。

この文書で説明されていないコードスタイルの側面や、自動的に強制されないものは、通常コードレビュー中にキャッチされるか、強制されません。

C++ガイドライン

一般に、C++ Core Guidelinesに従うことをお勧めします。また、Sean Parentの「C++ Seasoning」トークを見ることをお勧めし、そのルールに従うよう努めています。「生のループなし。生のポインタなし。生の同期プリミティブなし。」

STLおよびThrustのアルゴリズムを生のループよりも優先します。
生のポインタや生のメモリアロケーションよりもlibcudfおよびRMMの所有データ構造とビューを優先します。
libcudfには多くのCPUスレッドの並行性はありませんが、いくつかあります。現在、libcudfは生の同期プリミティブを使用しています。そのため、Parentの第三のルールを再検討し、改善する必要があります。

libcudfコードの追加スタイルガイドライン：

「イーストconst」を優先し、constを型の後に置きます。これはclang-formatによって自動的に強制されません。なぜなら、QualifierAlignment: Rightオプションが偽陽性と偽陰性を生成することが観察されているためです。
NL.11: リテラルを読みやすくする: 10進数値は桁数ごとに整数セパレータを使用する必要があります（例：1'234'567）。16進数値は4文字ごとにセパレータを使用する必要があります（例：0x0123'ABCD）。

ドキュメントに関する詳細は、Documentation Guideで説明されています。

インクルード

以下のガイドラインは、#include行の整理に適用されます。

ライブラリごとにインクルードをグループ化します（例：cuDF、RMM、Thrust、STL）。clang-formatはグループを尊重し、グループ内の個々のインクルードを辞書順にソートします。
グループごとに空行で区切ります。
グループを「最も近い」から「最も遠い」順に配置します。つまり、ローカルのインクルード、次に他のRAPIDSライブラリからのインクルード、次にのような関連ライブラリからのインクルード、次にcuDFと共にインストールされた依存関係からのインクルード、最後に標準ヘッダー（例：、）を含めます。
clang-formatを使用してヘッダーを自動的にグループ化およびソートします。詳細については、cudf/cpp/.clang-formatファイルを参照してください。
インクルードパスの最初の部分がlibcudfのincludeディレクトリにない限り、すべてのインクルードに<>を使用してください。言い換えれば、libcudfの内部ヘッダー（例：srcまたはtestディレクトリ内）であれば、パスはのようにcudfで始まらないため、クォートを使用する必要があります。例：#include "io/utilities/hostdevice_vector.hpp"。
cudf_testとnvtextはlibcudfリポジトリ内の別のライブラリであるため、publicヘッダーがincludeにあり、<>でインクルードする必要があります。
clangdのようなツールは、インクルードを自動的に挿入することがよくありますが、通常、グループ化とブラケットを間違えます。クォートまたはブラケットの使用を修正し、次にclang-formatを実行してグループ化を修正します。
インクルードが含まれるファイルに対してのみ必要であることを常に確認します。特にヘッダーファイルでの過剰なインクルードを避けてください。コードを削除するときにはこれを二重に確認してください。
可能な限り..を含む相対パスを避けます。..を含むパスは、同じディレクトリにないソースパスからのヘッダーをインクルードする場合に必要です。なぜなら、ソースパスはIで渡されないためです。
非内部ファイルからライブラリ内部ヘッダーをインクルードすることを避けてください。たとえば、libcudfのsrcディレクトリからのヘッダーをテストやlibcudfのpublicヘッダーにインクルードしないようにします。このようなことをする場合は、内部ヘッダーの一部をpublicヘッダーに移動することを検討してください。

libcudfデータ構造

libcudfのアプリケーションデータはカラムやテーブルに含まれていますが、libcudfコードを開発する際に使用する他のさまざまなデータ構造があります。

ビューと所有権

リソースの所有権は、libcudfの重要な概念です。簡単に言うと、「所有」オブジェクトはリソース（たとえばデバイスメモリ）を所有します。所有オブジェクトは、構築中にそのリソースを取得し、破棄中にリソースを解放します（RAII）。「非所有」オブジェクトはリソースを所有しません。libcudf内の*_viewサフィックスを持つクラスはすべて非所有です。詳細については、libcudfプレゼンテーションを参照してください。

libcudf関数は通常、入力としてビュー（column_viewまたはtable_view）を受け取り、出力として所有オブジェクトへのunique_ptrを生成します。たとえば、

cppコードをコピーする
std::unique_ptr<table> sort(table_view const& input);

メモリリソース

libcudfは、すべてのデバイスメモリをRMMメモリリソース（MR）またはCUDA MR経由で割り当てます。どちらのタイプも、rmm::device_async_resource_refパラメーターを介してlibcudf関数に渡すことができます。詳細については、RMMドキュメントを参照してください。

現在のデバイスメモリリソース

RMMは、デバイスごとに「デフォルト」メモリリソースを提供しており、rmm::mr::get_current_device_resource()およびrmm::mr::set_current_device_resource(...)関数を介してアクセスおよび更新できます。すべてのメモリリソースパラメーターは、rmm::mr::get_current_device_resource()の戻り値をデフォルトとして使用する必要があります。

リソース参照

メモリリソースは、リソース参照パラメーターを介して渡されます。リソース参照は、消費者が期待するリソースのプロパティを指定できるメモリリソースラッパーです。これらはlibcu++のcuda::mr名前空間で定義されていますが、RMMはrmm/resource_ref.hppにいくつかの便利なラッパーを提供しています。

rmm::device_resource_ref: デバイスアクセス可能なメモリの同期割り当てを提供するメモリリソースを受け入れます。
rmm::device_async_resource_ref: ストリーム順序でデバイスアクセス可能なメモリの割り当てを提供するメモリリソースを受け入れます。
rmm::host_resource_ref: ホストアクセス可能なメモリの同期割り当てを提供するメモリリソースを受け入れます。
rmm::host_async_resource_ref: ストリーム順序でホストアクセス可能なメモリの割り当てを提供するメモリリソースを受け入れます。
rmm::host_device_resource_ref: ホストおよびデバイスアクセス可能なメモリの同期割り当てを提供するメモリリソースを受け入れます。
rmm::host_async_resource_ref: ストリーム順序でホストおよびデバイスアクセス可能なメモリの割り当てを提供するメモリリソースを受け入れます。

リソース参照の詳細については、libcu++ドキュメントを参照してください。

cudf::column

cudf::columnは、libcudfの主要な所有データ構造です。ほとんどのlibcudf公開APIは、出力としてcudf::columnまたはcudf::tableを生成します。カラムは、カラムの要素のデバイスメモリを所有するdevice_buffersと、オプションでヌルインジケータビットマスクを含みます。

column_viewおよびmutable_column_viewに暗黙的に変換可能です。

移動可能でコピー可能です。コピーはカラムの内容のディープコピーを実行し、ムーブはカラムから別のカラムに内容を移動します。

例：

cppコードをコピーする
cudf::column col{...};

cudf::column copy{col}; // `col`の内容をコピー
cudf::column const moved_to{std::move(col)}; // `col`から内容をムーブ

column_view v = moved_to; // 非所有の`column_view`への暗黙的な変換
// `mutable_column_view m = moved_to;` // constカラムへのmutableビューの作成は不可

カラムは、データ型に応じてネストされた（子）カラムを持つ場合があります。たとえば、リスト、構造体、文字列型カラムがあります。

cudf::column_view

cudf::column_viewは、libcudfの主要な非所有データ構造です。これは、カラムのデバイスメモリの不変の非所有ビューです。ほとんどのlibcudf公開APIは、入力としてビューを受け取ります。

column_viewはカラムの「スライス」のビューである可能性があります。たとえば、1000行のカラムのうち75〜150行目のビューである場合があります。このcolumn_viewのサイズは75であり、ビューのインデックス0をアクセスすると、所有カラムのインデックス75の要素が返されます。内部的には、ビューにポインタ、オフセット、サイズが保存されて実装されています。column_view::data<T>()はcolumn_view::head<T>() + offsetへのポインタイテレータを返します。

cudf::mutable_column_view

カラムのデバイスメモリのmutableな非所有ビューです。詳細APIおよび（稀な）カラムをその場で変更する公開APIに使用されます。

cudf::column_device_view

要素のカラムとしてデバイスコードで使用可能な、カラムデータの不変の非所有ビューです。CUDAカーネルやデバイス関数（Thrustアルゴリズムを含む）への入力としてcolumn_viewデータを渡すために使用されます。

cudf::mutable_column_device_view

要素のカラムとしてデバイスコードで使用可能な、カラムデータのmutableな非所有ビューです。CUDAカーネルやデバイス関数（Thrustアルゴリズムを含む）でcolumn_viewデータを変更するために使用されます。

cudf::table

同じ数の要素を持つcudf::columnのセットを所有するクラスです。これは、データフレームに相当するC++オブジェクトです。

cudf::table_viewおよびcudf::mutable_table_viewに暗黙的に変換可能です。

移動可能でコピー可能です。コピーはすべてのカラムのディープコピーを実行し、ムーブはすべてのカラムを1つのテーブルから別のテーブルに移動します。

cudf::table_view

テーブルの不変の非所有ビューです。

cudf::mutable_table_view

テーブルのmutableな非所有ビューです。

cudf::size_type

cudf::size_typeは、カラム内の要素数、要素のオフセット、特定の要素をアドレスするインデックス、カラム要素のサブセットのセグメントなどに使用される型です。これは32ビットの符号付き整数型に相当し、最大値は2147483647です。いくつかのAPIは負のインデックス値も受け入れ、その関数は-2147483648の最小値をサポートします。この基本的な型は、カラムサイズの制限だけでなく、要素をカウントするための出力値にも影響を与えます。

スパン

libcudfは、C++20のstd::spanに似たスパンクラスを提供しており、オブジェクトの連続したシーケンスの軽量ビューです。libcudfは、host_spanおよびdevice_spanの2つのクラスを提供しており、複数のコンテナ型から、またはポインタ（それぞれホストまたはデバイス）とサイズ、またはイテレータから構築できます。スパン型は、複数の入力コンテナ型で動作する汎用（内部）インターフェースを定義するのに便利です。device_spanはthrust::device_vector、rmm::device_vector、またはrmm::device_uvectorから構築できます。host_spanはthrust::host_vector、std::vector、またはstd::basic_stringから構築できます。

libcudfコードでスパンを使用する内部（詳細）関数を定義する場合、関数をより広範に適用可能にするために、特定のベクタ型の代わりに入力ベクタパラメータにスパンを使用してください。

スパンが不変の要素を参照している場合、スパン自体ではなく、テンプレート型パラメータにconstを適用します。また、スパンは軽量ビューであるため、値で渡すべきです。libcudfでスパンを入力として受け取るAPIは、次のような非同期でデバイスデータをホストのstd::vectorにコピーする関数になります。

template <typename T>
std::vector<T> make_std_vector_async(device_span<T const> v, rmm::cuda_stream_view stream)

cudf::scalar

cudf::scalarは、cudfで現在サポートされている任意の型の単一の、nullableな値を表すオブジェクトです。各値の型は、それぞれ異なるスカラークラスで表され、すべてcudf::scalarから派生しています。例：numeric_scalarは単一の数値を保持し、string_scalarは単一の文字列を保持します。保存される値のデータはデバイスメモリにあります。

list_scalarは、単一のリストの基礎データを保持します。つまり、基礎データはcudfがサポートする任意の型になり得ます。たとえば、整数のリストを表すlist_scalarは、INT32型のcudf::columnを保持します。整数のリストのリストを表すlist_scalarは、LIST型のcudf::columnを保持し、次にINT32型のカラムを保持します。

コンストラクション

スカラーは、それぞれのコンストラクターを使用して、またはmake_numeric_scalar()、make_timestamp_scalar()、make_string_scalar()などのファクトリーメソッドを使用して作成できます。

キャスティング

すべてのファクトリーメソッドは、unique_ptr<scalar>を返し、その値にアクセスする前にそれぞれのスカラクラス型に静的ダウンキャストする必要があります。キャスティングなしで有効性（nullness）にアクセスできます。一般的に、値は値型を認識している関数からアクセスする必要があります。たとえば、type_dispatcherからディスパッチされたファンクターなどです。値型が与えられたときに必要なスカラクラス型にキャストするには、type_dispatcher.hppで提供されるマッピングユーティリティscalar_type_tを使用します。

// unique_ptr<scalar> s = make_numeric_scalar(...);

using ScalarType = cudf::scalar_type_t<T>;
// ScalarTypeは現在numeric_scalar<T>です。
auto s1 = static_cast<ScalarType *>(s.get());

デバイスへのパス

list_scalarを除くすべてのスカラ型には、デバイスから値と有効性にアクセスするための非所有デバイスビュークラスがあります。これは、get_scalar_device_view(ScalarType s)関数を使用して取得できます。デバイスビューは基本的なスカラーオブジェクトには提供されず、派生した型付きスカラクラスオブジェクトにのみ提供されます。

list_scalarの基礎データは、view()メソッドを介してアクセスできます。非ネストデータの場合、デバイスビューはcolumn_device_view::create(column_view)関数を使用して取得できます。ネストデータの場合、リストカラム用の特殊なデバイスビューはlists_column_device_view(column_device_view)を介して構築できます。

libcudfのポリシーと設計原則

libcudfは、データサイエンスで発生するさまざまな問題を解決するための、スレッドセーフで単一GPU加速されたアルゴリズムのプリミティブを提供するように設計されています。APIは、デフォルトのGPUで実行されるように記述されており、これは標準のCUDAデバイスAPIまたは環境変数CUDA_VISIBLE_DEVICESを通じて呼び出し元が制御できます。私たちの目標は、SparkやPandasのような多様なユースケースがGPUのパフォーマンスを活用できるようにすることです。そして、libcudfは、これらの高レベルレイヤー（SparkやDaskなど）がマルチGPUタスクを調整するのに依存しています。

これらのユースケースを最適に満たすために、libcudfはパフォーマンスと柔軟性を優先しており、時には利便性を犠牲にすることもあります。私たちはユーザーがlibcudfを直接使用することを歓迎しますが、ほとんどのユーザーが、SparkやcuDF Pythonのような高レベルレイヤーを通じてlibcudfを消費し、libcudfを直接使用するユーザーが自分で対処しなければならない詳細の一部を処理することを期待して設計しています。これらのポリシーとその理由をここに文書化します。

libcudfはデータを内省しない

libcudf APIは通常、入力データの深い内省や検証を行いません。これにはいくつかの理由があります：

単一責任の原則に違反するためです：検証は実行とは別です。
libcudfのデータ構造はデータをGPUに保存しているため、検証には最低限カーネルの起動のオーバーヘッドが発生し、一般的には禁止されるほど高価になる可能性があります。
データの内省に関するAPIの約束は、実装を大幅に複雑にすることがよくあります。

したがって、ユーザーはこれらのAPIに対して有効なデータを渡す責任を負います。このポリシーはlibcudfがまったく検証を行わないことを意味するわけではありません。libcudf APIは、内省を必要としない検証は引き続き行うべきです。libcudfが検証するべきことと、検証するべきでないことの例をいくつか挙げます。

libcudfが検証するべきこと：

入力カラム/テーブルのサイズまたはデータ型

libcudfが検証しないべきこと：

整数のオーバーフロー
出力が特定の入力セットに対して2GBのサイズ制限を超えないことの保証

libcudfはネストされた型のヌルマスクをサニタイズすることを期待する

ネストされたデータ型のカラム（LISTやSTRUCTなど）を受け入れるさまざまなlibcudf APIは、これらのカラムがサニタイズされていると仮定する場合があります。この文脈でのサニタイズとは、ネストされたカラムのヌル要素がネストされたカラムの要素と互換性があることを確認することを指します。具体的には：

リストカラムのヌル要素は空であるべきです。ヌル要素の開始オフセットは終了オフセットと等しいべきです。
構造体カラムのヌル要素は、基礎となる構造体のヌル要素であるべきです。
複合カラムの場合、ヌルは親カラムのレベルでのみ存在するべきです。子カラムにはヌルが含まれてはなりません。

ネストされたカラムに対するスライス操作は、子カラムにオフセットを伝播しません。

libcudf APIは「汚れた」カラム（つまり、サニタイズされていないデータを含むカラム）を返さないことを約束するべきです。したがって、問題は、ユーザーがサニタイズされていない入力カラムを構築し、それをlibcudf APIに渡す場合のみです。

libcudf APIを非同期であるかのように扱う

ホストで呼び出されたlibcudf APIは、ストリームが戻る前に同期されることを保証しません。libcudf内での作業はcudf::get_default_stream().valueで行われ、これはデフォルトでCUDAデフォルトストリーム（ストリーム0）です。ストリーム0の動作は、CUDF_USE_PER_THREAD_DEFAULT_STREAMを使用してスレッドごとのデフォルトストリームが有効化されている場合に異なります。非ブロッキングストリームを使用しているlibcudfによって提供されるデータまたはlibcudfによって返されるデータには、ストリームの安全性を確保するためにlibcudfのデフォルトストリームとの同期が必要です。

libcudfは通常、特定の順序付けを保証しない

libcudf内のmergeやgroupbyのような関数は、出力内のエントリの順序に関して保証を行いません。決定論的な順序を約束することは、一般的には高速な並列アルゴリズムに適していません。出力がソートされる必要がある場合は、呼び出し元が事後にソートを実行する責任があります。

libcudfは特定の例外メッセージを約束しない

libcudfは、さまざまな種類の無効な入力に対してAPIがスローする例外を文書化しています。これらの例外の型（たとえばcudf::logic_errorなど）は公開APIの一部です。ただし、これらの例外のwhatメソッドによって返される説明文字列はAPIの一部ではなく、変更される可能性があります。libcudfのエラーメッセージの内容に依存してエラーの性質を判断することは避けてください。libcudfがさまざまな状況でスローする例外の種類については、エラーハンドリングのセクションを参照してください。

libcudf APIと実装

ストリーム

libcudfは現在、CUDAストリームを使用した非同期実行のサポートを追加しています。ストリームの使用を促進するために、デバイスメモリを割り当てたりカーネルを実行したりするすべての新しいlibcudf APIは、rmm::cuda_stream_viewパラメーターを最後に受け取り、そのデフォルト値をcudf::get_default_stream()に設定する必要があります。このルールには1つの例外があります：APIがメモリリソースパラメーターも受け入れる場合、ストリームパラメーターはメモリリソースの直前に配置する必要があります。このAPIは次に、同じ署名を持つ対応する詳細APIに呼び出しを転送する必要がありますが、詳細APIにはストリームのデフォルトパラメーターが含まれていてはなりません（詳細APIは常にデフォルトパラメーターを避けるべきです）。実装は、詳細API定義内に完全に含まれており、ストリームパラメーターを使用して非同期バージョンのCUDA APIのみを使用する必要があります。

他のlibcudf関数から詳細APIを呼び出すことができるようにするために、その詳細APIをcudf/cpp/include/detail/ディレクトリに配置されたヘッダーで公開する必要があります。他のlibcudf関数が詳細関数を呼び出さない場合は、宣言は必要ありません。

例えば：

// cpp/include/cudf/header.hpp
void external_function(...,
  rmm::cuda_stream_view stream      = cudf::get_default_stream(),
  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());

// cpp/include/cudf/detail/header.hpp
namespace detail{
void external_function(..., rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr)
} // namespace detail

// cudf/src/implementation.cpp
namespace detail{
// ストリームパラメーターを使用して、詳細実装を行います。
void external_function(..., rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr){
  // 実装は、ストリームを使用して非同期APIを使用します。
  rmm::device_buffer buff(..., stream, mr);
  CUDF_CUDA_TRY(cudaMemcpyAsync(...,stream.value()));
  kernel<<<..., stream>>>(...);
  thrust::algorithm(rmm::exec_policy(stream), ...);
}
} // namespace detail

void external_function(..., rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr)
{
  CUDF_FUNC_RANGE(); // この関数のライフタイムのためにNVTX範囲を生成します。
  detail::external_function(..., stream, mr);
}

注：非同期デバイスからホストへのコピーの結果として非ポインタ値がAPIから返される場合、ストリームを同期してから戻る必要があります。ただし、カラムが返される場合、ストリームを同期しないでください。同期することで非同期性が失われます。

注：cudaDeviceSynchronize()は絶対に使用しないでください。これにより、libcudf APIを使用するマルチストリーム/マルチスレッド作業の能力が制限されます。

ストリームの作成

libcudf機能を実装する際に、アルゴリズムを実装する際にオーバーラップを達成するために内部的にストリームを使用することが有利な場合があります。ただし、ストリームを動的に作成することはコストがかかる場合があります。RMMには、動的なストリーム作成を回避するためのストリームプールクラスがありますが、これらはlibcudfでまだ公開されていません。そのため、現在のところ、libcudf機能はストリームを作成することを避けるべきです（効率が少し劣る場合でも）。ストリームを使用するのが有益な場所には、// TODO:のコメントを残すのが良い考えです。

メモリアロケーション

libcudfは、デバイスメモリアロケーションの方法を抽象化し、制御するためにデバイスメモリリソースを使用します。

出力メモリ

ユーザーに返されるメモリを割り当てる任意のlibcudf APIは、最後にrmm::device_async_resource_refを受け入れる必要があります。このAPI内では、このメモリリソースを使用して返されるオブジェクトのメモリを割り当てる必要があります。そのため、返されるオブジェクトのメモリを割り当てる関数に渡す必要があります。例：

// 返された`column`には新しく割り当てられたメモリが含まれているため、APIはメモリリソースポインタを受け入れる必要があります。
std::unique_ptr<column> returns_output_memory(
  ..., rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());

// このAPIは新しい*出力*メモリを割り当てないため、メモリリソースは不要です。
void does_not_allocate_output_memory(...);

このルールは、メモリリソースを割り当てるすべての詳細APIに自動的に適用されます。任意の詳細APIは任意の公開APIによって呼び出される可能性があり、したがってユーザーに返されるメモリリソースを割り当てる可能性があります。このようなユースケースをサポートするために、メモリリソースを割り当てるすべての詳細APIはmrパラメーターを受け入れるべきです。呼び出し元は、必要に応じて、提供されたmrまたはrmm::mr::get_current_device_resource()を介して渡す責任があります。

一時メモリ

libcudf API内で割り当てられるすべてのメモリが呼び出し元に返されるわけではありません。アルゴリズムは、しばしば中間結果のための一時的なスクラッチメモリを割り当てる必要があります。常に、rmm::mr::get_current_device_resource()から取得したデフォルトリソースを一時メモリアロケーションに使用してください。例：

rmm::device_buffer some_function(
  ..., rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()) {
    rmm::device_buffer returned_buffer(..., mr); // 返されたバッファは、渡されたMRを使用します
    ...
    rmm::device_buffer temporary_buffer(...); // 一時バッファはデフォルトMRを使用します
    ...
    return returned_buffer;
}

メモリ管理

libcudfコードでは、生ポインタや直接的なメモリアロケーションを避けます。デバイスメモリアロケーションのためにメモリリソースを使用し、自動ライフタイム管理を行うRMMクラスを使用します。

rmm::device_buffer

指定されたバイト数の未初期化デバイスメモリを、メモリリソースを使用して割り当てます。rmm::device_async_resource_refが明示的に提供されていない場合、rmm::mr::get_current_device_resource()を使用します。

rmm::device_bufferは、ストリーム上で移動可能でコピー可能です。コピーはdevice_bufferのデバイスメモリのディープコピーを指定されたストリーム上で実行し、移動はデバイスメモリの所有権を一方のdevice_bufferから他方へ移動します。

// 指定されたリソースとストリームを使用して、少なくとも100バイトの未初期化デバイスメモリを割り当てます
rmm::device_buffer buff(100, stream, mr);
void * raw_data = buff.data(); // 基礎デバイスメモリへの生ポインタ

// `buff`を`copy`にディープコピーします。
rmm::device_buffer copy(buff, stream);

// `buff`の内容を`moved_to`に移動します。
rmm::device_buffer moved_to(std::move(buff));

custom_memory_resource *mr...;
// カスタムメモリリソースから100バイトを割り当てます
rmm::device_buffer custom_buff(100, mr, stream);

rmm::device_scalar

指定された値に初期化された指定された型の単一要素を割り当てます。これは、デバイスカーネルへのスカラー入力/出力のために使用されます。たとえば、リダクション結果やヌルカウントなどです。これは、長さ1のrmm::device_vector<T>の便利なラッパーです。

// 指定されたリソースとストリームを使用して、1つのintのデバイスメモリを割り当て、値を42に初期化します
rmm::device_scalar<int> int_scalar{42, stream, mr};

// scalar.data()はデバイスメモリ内の値へのポインタを返します
kernel<<<...>>>(int_scalar.data(),...);

// scalar.value()は、スカラーのストリームを同期し、値をデバイスからホストにコピーして値を返します
int host_value = int_scalar.value();

rmm::device_vector

指定された数の指定された型の要素を割り当てます。初期化値が提供されていない場合、すべての要素はデフォルトで初期化されます（これにはカーネル起動が含まれます）。

注：libcudfからrmm::device_vectorおよびthrust::device_vectorのすべての使用を削除しており、libcudfの新しいコードで使用する際には慎重に検討してください。代わりに、rmm::device_uvectorとdevice_factories.hppのユーティリティファクトリーを使用してください。これらのユーティリティは、ホスト側ベクターからのuvectorの作成、またはゼロ初期化されたuvectorの作成を可能にし、device_vectorと同じくらい便利に使用できます。device_vectorを避けることには、次のrmm::device_uvectorのセクションで説明されるように、いくつかの利点があります。

rmm::device_uvector

device_vectorに似ており、デバイスメモリ内で一連の要素を連続して割り当てますが、いくつかの重要な違いがあります。

最適化として、要素は初期化されず、構築時には同期が発生しません。これにより、型Tはtrivially copyable型に制限されます。
すべての操作はストリーム順序で行われます（つまり、操作が実行されるストリームを指定するcuda_stream_viewを受け取ります）。これにより、非デフォルトストリームを使用する際の安全性が向上します。
device_uvector.hppには、__device__コードが含まれていません。これにより、thrust/device_vector.hppとは異なり、device_uvectorを.cuファイルだけでなく.cppファイルでも使用できます。

cuda_stream s;
// ストリーム`s`上で`int32_t`要素100個の未初期化ストレージをデフォルトリソースを使用して割り当てます
rmm::device_uvector<int32_t> v(100, s);
// 要素を0に初期化します
thrust::uninitialized_fill(thrust::cuda::par.on(s.value()), v.begin(), v.end(), int32_t{0});

rmm::mr::device_memory_resource * mr = new my_custom_resource{...};
// ストリーム`s`上で`mr`リソースを使用して`int32_t`要素100個の未初期化ストレージを割り当てます
rmm::device_uvector<int32_t> v2{100, s, mr};

デフォルトパラメーター

libcudfの公開APIにはデフォルトの関数パラメーターを含めることができますが、詳細関数には含めるべきではありません。デフォルトのメモリリソースパラメーターは、開発者が間違ったリソースを使用してメモリを誤って割り当てるのを容易にするため、詳細APIではデフォルトのメモリリソースを避け、すべてのメモリアロケーションを慎重に検討する必要があります。

現在、libcudfのAPIでストリームは公開されていませんが、将来的には公開される予定です。その結果、メモリリソースにも適用されるのと同じ理由で、ストリームも適用されます。公開APIではcudf::get_default_stream()を使用するようにデフォルト設定されています。ただし、詳細APIに同じデフォルト設定を含めると、ユーザー提供のストリームを渡すことを忘れる可能性があります。すべての詳細API呼び出しで明示的にストリームを渡すようにすることは、そのような間違いを防ぐことを目的としています。

メモリリソース（および最終的にはストリーム）は、ほぼすべての公開APIの最後のパラメーターです。APIの一貫性のため、libcudfの内部でも同じです。したがって、デフォルトストリームまたはMRを許可しないことの結果として、詳細APIのパラメーターにはデフォルトがありません。

You get articles that match your needs
You can efficiently read back useful information
You can use dark theme

What you can do with signing up