『Python NumbaとCUDA Cを使用したバッチK-Means』

『Python NumbaとCUDA Cを活用したバッチK-Means』

データ解析の高速化方法：1600倍 vs Scikit-Learn（コード付き！）

データ解析のワークロードを並列化することは、特定のユースケースに対して効率的な既製の実装がない場合には困難な作業です。このチュートリアルでは、C言語とPython Numbaの両方でCUDAカーネルを作成する原則と、それらの原則をクラシックなk-meansクラスタリングアルゴリズムに適用する方法について説明します。この記事の最後まで読むと、標準のscikit-learn実装と比較して、CとPythonの両方でバッチ処理されたk-meansのカスタム並列化実装を書くことができ、最大で1600倍の高速化を達成することができます。コードにすぐにアクセスしたい場合は、Colabで利用可能です。

イントロダクション

NVIDIA GPUの並列化フレームワークとして2つのフレームワークを学びます：Python NumbaとCUDA C。それぞれを実装し、scikit-learnとのベンチマークを取ります。

Numbaは、CよりもPythonを好む人々にとって学習曲線がなめらかです。Numbaは、コードの一部を特殊化されたCUDA関数であるカーネルにコンパイルします。CUDA Cは抽象化ツリーのより低いレイヤーであり、より細かな制御を提供します。Colabの例はアルゴリズムの実装を密接に反映しているように設計されているため、どちらかを理解すればもう一方を簡単に理解することができます。

このチュートリアルは、私が書かなければならなかったカスタムバッチk-meansの実装から派生しています。通常、k-meansライブラリは非常に大規模なデータセットでアルゴリズムの単一のインスタンスを実行するように最適化されています。しかし、私たちのプロジェクトでは、並列に数百万の個々の小さなデータセットをクラスタリングする必要があり、それに対して既製のライブラリを見つけることはできませんでした。

• LMQL — 言語モデル用のSQL
• 「MLを学ぶ勇気：L1とL2の正則化の解明（パート1）」
• あなたのビジネス分析を高めましょう：季節調整のステップバイステップガイド

この記事は、CUDAの基礎、並列化されたK-meansアルゴリズム、Python Numba、CUDA Cの4つのメインセクションに分かれています。私は素材を比較的自己完結的にして、必要に応じて概念を簡単に再度説明します。

CUDAの基礎

CPUは高速な逐次処理に設計されており、GPUは大量の並列処理に設計されています。私たちのアルゴリズムでは、数百万の小さな独立したデータセットでk-meansを実行する必要があり、これはGPUの実装に非常に適しています。

私たちの実装では、CUDAという名前の、NVIDIAによって開発されたCライブラリとコンピューティングプラットフォームを使用します。これは、GPUを使用した並列計算を活用するために開発されたものです。

GPUカーネルとデバイス関数を理解するために、以下の定義を確認しましょう：

スレッドは、独立して実行できる単一の命令のシーケンスです。

コアは、単一のスレッドを実行できる処理ユニットです。

ワープは、最小のスレッドスケジューリング単位です。各ワープは32個のスレッドから成り、それらをコアにスケジューリングします。

マルチプロセッサまたはStreaming Multiprocessor（SM）は、固定数のコアで構成されています。

GPUカーネルは、多くのGPUスレッドごとに並列に実行されるように設計された関数です。カーネル関数はCPUによって呼び出され、GPU上で実行されます。

デバイス関数は、カーネル関数または他のデバイス関数から呼び出すことができる関数です。この関数はGPU上で実行されます。

これらのカーネルは、グリッド、ブロック、およびスレッドの階層に組織されています。GPUのコンテキストでは、各スレッドは単一のコアに接続され、カーネルのコピーを実行します。

ブロックは、1つのマルチプロセッサで実行されるスレッドの集合です。

グリッドは、スレッドのブロックを組織化するための抽象配列です。グリッドは、インデックスによってカーネルインスタンスをスレッドにマップするために使用されます。

GPUには、グローバルメモリと共有メモリの2つのタイプのメモリがあります。

グローバルメモリはDRAM（動的ランダムアクセスメモリ）に格納されています。すべてのスレッドはその内容にアクセスできますが、メモリのレイテンシが高いです。

共有メモリは、キャッシュに格納され、同じブロック内のスレッドに対してプライベートです。グローバルメモリよりもはるかに高速にアクセスできます。

並列化されたK-meansアルゴリズム

では、k-meansアルゴリズムを紹介しましょう。 k-meansは、データセットをk個の異なる、重なり合わないクラスタに分割する非監視学習のアルゴリズムです。データポイントのセットが与えられた場合、最初にk個のセントロイド、または開始クラスタの中心を初期化します：

その後、初期セントロイドを選択した後、次の2つのステップを反復的に実行します：

1. 割り当てステップ：ユークリッド距離に基づいて、各データポイントを最も近いセントロイドに割り当てます。
2. 更新ステップ：前のステップでそのセントロイドに割り当てられたすべてのポイントの平均値にセントロイドの位置を再割り当てします。

これらのステップは収束するまで繰り返されます。つまり、セントロイドの位置が大幅に変化しなくなったときです。出力は、k個のクラスタセントロイドの座標のセットと、元のデータポイントごとにクラスタインデックスをマークする配列（ラベルと呼ばれる）です。

大規模なデータセットでは、セントロイドの初期化はアルゴリズムの出力に大きな影響を与える場合があります。したがって、プログラムは複数の初期セントロイド、つまり初期のシードを試し、最良のシードの結果を返します。各シードは、初期データセットから非選出で選択されます – つまり、初期セントロイドは繰り返されません。アルゴリズムに最適なシードの数は、データポイントの数の1/3です。プログラムでは、100個のデータポイントの個々の行でk-meansを実行するため、最適なシードの数は33になります。

私たちのk-means関数では、100万行はブロックで表され、スレッドはシードを表します。ブロック内のスレッドはワープに組織化され、ハードウェアアーキテクチャの最小スレッドスケジュール単位です。各ワープは32のスレッドで構成されており、ブロックサイズを32の倍数に設定するのが最適です。それぞれのブロックは、クラスタセンタとそれに割り当てられたポイントのユークリッド距離の合計によって測定される最小慣性を持つシードのデータを出力します。

こちらには、PythonまたはCのいずれかで進めることができるColabへのリンクがあります。

global variables: numRows, lineSize, numClustersdef hostKMeans:    inputData = initializeInputData(numRows, lineSize)    outputCentroids = createEmptyArray(numRows, numClusters)    outputLabels = createEmptyArray(numRows, lineSize)        sendToDevice(inputData, outputCentroids, outputLabels)    cuda_kmeans[blockDimensions, threadDimensions](inputData, outputCentroids, outputLabels)    waitForKernelCompletion()    copyFromDevice(outputCentroids, outputLabels)

私たちのk-meansアルゴリズムでは、まずグローバル変数を設定します。CPUとGPUの両方から参照する必要があります。

グローバル変数はカーネルからアクセスできますが、カーネルは直接ホストに値を返すことはできません。この制限を回避するために、コードのCPU部分は、入力データとともにカーネルに2つの空の配列を渡します。これらの配列は、最終的なセントロイドとラベルの配列をCPUにコピーするために使用されます。

データ構造がインスタンス化されたら、グリッドとブロックの次元をタプルとして定義してカーネルを呼び出します。カーネルへの呼び出しには、データの渡し込み、セントロイドとラベルのメモリ割り当て、および初期ランダムセントロイド初期化のための状態が含まれます。

def KMeansKernel(data, outputCentroids, outputLabels)    
    row = currentBlock()    
    seed = currentThread()    
    sharedInputRow = sharedArray(shape=(lineSize))    
    sharedInertia = sharedArray(shape=(numSeeds))    
    sharedCentroids = sharedArray(shape=(numSeeds, numClusters))    
    sharedLabels = sharedArray(shape=(numSeeds, lineSize))    
    sharedInputRow = data[row]    
    synchronizeThreads()    
    if seed == 0        
        centroids = initializeCentroids(data)    
    synchronizeThreads()        
    KMeansAlgorithm(sharedInputRow, sharedCentroids, sharedLabels)        
    sharedInertia[Seed] = calculateInertia(sharedInputRow, sharedCentroids, sharedLabels)        
    synchronizeThreads()    
    if seed == 0        
        minInertiaIndex = findMin(Inertia)    
    sharedOutputCentroids = centroids[minInertiaIndex]    
    sharedOutputLabels = labels[minInertiaIndex]

デバイス（GPU）上にいると、コードは今やグリッド全体で同時に存在しています。グリッド上の位置を特定するために、ブロックとスレッドのインデックスにアクセスします。

GPUでは、メモリはデフォルトでグローバルメモリ（DRAMに格納）になります。 共有メモリは同じブロック内のスレッドに対してプライベートです。

全てのスレッドが読み取る必要がある単一のデータ行を共有メモリに転送することで、メモリアクセス時間を削減します。cuda_kmeans関数では、データ行、セントロイド、ラベル、および各シードの精度測定である慣性を格納するために共有メモリを作成します。

プログラムでは、各スレッドが1つのシードに対応し、ブロック内のすべてのスレッドが同じデータ行で作業します。各ブロックでは、1つのスレッドが32個のシードを順次作成し、その結果をブロック内の他のスレッドのための単一のデータ構造に集約します。

アルゴリズムの次のステップがこの集約が完了していることに依存する場合、スレッドは組み込みのCUDA syncthreads（）関数を使用して同期する必要があります。 NB:すべてのスレッドが完了する前にスレッドを同期しようとすると、デッドロックやプログラム全体のハングアップを引き起こす可能性があるため、syncthreads（）呼び出しの配置には非常に注意する必要があります。

以下に疑似コードで示されたカーネル関数cuda_kmeansがあります。この関数は、上記で説明したプロセスを整理し、すべての32個シードの結果を収容するスペースを作り、最終的な出力のセントロイドとラベルのための最適なシードを選択します。

def KMeansDevice(dataRow, centroids, labels)    
    seed = currentThread()    
    centroidsRow = centroids[seed]    
    labelsRow = labels[seed]     
    centroidsRow = sort(centroidsRow)    
    yardStick = computeYardstick(sortedCentroids)     
    oldCentroids = localArray(shape=(numSeeds, numClusters))    
    for iteration in range(100):        
        if converged(oldCentroids, centroidsRow)            
            break        
        oldCentroids = copy(centroidsRow)        
        assignLabels(dataRow, centroidsRow, labelsRow)        
        updateCentroids(dataRow, centroidsRow, labelsRow)

cuda_kmeansから実際のk-meansアルゴリズムを呼び出し、新しくインスタンス化された共有メモリを渡します。k-meansのアルゴリズムでは、最初のセントロイドを選択してから、それらを最小から最大まで並べ替えます。データポイントを最も近いセントロイドに反復的に割り当て、収束するまでセントロイドの位置を更新します。

収束が達成されたかどうかを判断するために、find_yard_stickというヘルパー関数を使用します。この関数は、2つの初期セントロイド（yard_stick）間の最小距離を計算して返します。収束条件は、イテレーション中にセントロイドのいずれもがyard_stick倍のイプシロンよりも多く移動していない場合に満たされます。

収束後、cuda_kmeansに戻ります。ここでは、各セントロイドとそのデータポイント間の二乗ユークリッド距離を計算して最適なシードを決定します。最も効果的なグループ化を示す、つまり最小の慣性を持つシードが最良と見なされます。次に、これらのセントロイドとラベルを取り出して、出力配列の単一行にコピーします。すべてのブロックが完了したら、これらの出力をホスト（CPU）にコピーします。

Numbaの紹介

カスタムカーネルの設計にはNumbaが最も簡単です。Numbaは、PythonのコードをCUDAカーネルにコンパイルするために使用できるPythonライブラリです。

NumbaはCUDAと連携しています。コードを並列化するために、Numbaは指定したGPUコードをカーネルにコンパイルし、それをGPUに渡します。これにより、プログラムの論理を次の2つのメインパーツに分割します：

1. CPUレベルのコード
2. GPUレベルのコード

Numbaを使用すると、コードの逐次実行および並列化可能な部分をCPUとGPUに分離し、それぞれに引き渡すことができます。GPUのために関数をコンパイルするには、プログラマは関数定義の上に@cuda.jitデコレータを使用し、この関数をCPU（ホスト）から呼び出されるが、GPU（デバイス）上で並列に実行されるカーネルに変換します。

Python Numba

Colabへのリンク。

NumbaはPythonコードとCUDAプラットフォームの橋渡し役となります。Pythonコードは、上記のアルゴリズムの疑似コードとほぼ同じですので、キーワードとなる構文の例をいくつか提供します。

cuda_kmeans[(NUM_ROWS,), (NUM_SEEDS,)](input_rows, output_labels, output_centroids, random_states)

必要なグローバル変数とデータ構造をインスタンス化した後、ホストからカーネルcuda_kmeansを呼び出すことができます。Numbaは、ブロックとスレッドの次元のために2つのタプルを要求します。1次元のブロックとスレッドを使用するため、各タプルの2番目のインデックスは空です。また、データ構造とランダムシードの配列も渡します。

@cuda.jit()def cuda_kmeans(input, output_labels, output_centroids, random_states):    row = cuda.blockIdx.x    seed = cuda.threadIdx.x    shared_input_row = cuda.shared.array(shape=(LINE_SIZE), dtype=np.float32)    shared_inertia = cuda.shared.array(shape=(NUM_SEEDS), dtype=np.float32)    shared_centroids = cuda.shared.array(shape=(NUM_SEEDS, NUM_CLUSTERS), dtype=np.float32)    shared_labels = cuda.shared.array(shape=(NUM_SEEDS, LINE_SIZE), dtype=np.int32)    if seed == 0:        get_initial_centroids(shared_input_row, shared_centroids, random_states)    cuda.syncthreads()    ...    kmeans(shared_input_row, shared_labels, shared_centroids)

GPUコンパイルをマークするために、Numbaのデコレータ@cuda.jit()を使用します。cuda.blockIdx.xおよびcuda.threadIdx.xの記法を使用して、カーネルの現在のインデックスを取得します。共有配列は、2つの引数（形状と型）を使用してcuda.shared.arrayを使用してインスタンス化されます。形状と型は、コンパイル時に既知である必要があります。各行にセントロイドを取得し、データで行を埋めた後、kmeans関数を呼び出し、共有配列を埋め、cuda.syncthreads()を呼び出します。

@cuda.jit(device=True)def kmeans(data_row, output_labels, output_centroids):     seed = cuda.threadIdx.x    labels_row = output_labels[seed]    centroids_row = output_centroids[seed]        ...    old_centroids = cuda.local.array(shape=(NUM_CLUSTERS), dtype=np.float32)    for iteration in range(NUM_ITERATIONS):            if iteration > 0:                if converged(centroids_row, old_centroids, yard_stick * EPSILON_PERCENT, iteration):                    break      # ラベルを割り当ててセントロイドを更新

k-meansはカーネル関数から呼び出されるため、デバイス関数です。したがって、デコレータにdevice=Trueを指定する必要があります：@cuda.jit(device=True)。k-means関数はラベルとセントロイドのための現在の行を取得し、収束するまで実行されます。

わずかな追加のコードと少しの努力で、あなたのPythonコードは最適化されたカーネルになり、並列使用に対応する準備ができます。

私たちの並列化されたk-meansは計算時間を大幅に削減しますが、Pythonのような高水準言語をラップしてコンパイルすることは必ずしも最適ではありません。プロジェクトのスピードアップのためにCでコードを書くことができるかどうか確かめるために、CUDA Cの世界に飛び込んでみました。

C入門

Pythonでは、メモリの割り当てと型の割り当ては自動で行われますが、Cではメモリをスタックまたはヒープに割り当てることができ、どちらも明示的な型宣言が必要で、固定量のメモリが割り当てられます。スタックメモリはコンパイラによって自動的に割り当てられ解放されますが、ヒープメモリはmalloc()などの関数を使ってランタイム時に手動で割り当てられ、解放はプログラマの責任です。

ポインタは変数のメモリアドレスを保持するツールです。ポインタが参照するデータの型は宣言時に定義されます。アスタリスク(*)を使用してポインタを指定します。変数のアドレスを取得するには、参照と呼ばれる操作でアンパサンド(&)を使用します。ポインタから値にアクセスするには、再度アスタリスクを使用してポインタを逆参照します。

二重ポインタ、つまりポインタのアドレスを格納するポインタは、他のポインタのアドレスを変更するのに役立ちます。これは、関数に配列を渡すときにアドレスを変更するためです。関数に配列を渡すとき、配列はサイズ情報を持たず、配列をインデックスで参照するためにポインタの算術演算に依存します。関数から配列を返すには、元のポインタのアドレスを渡すために&を使用し、ダブルポインタ**で受け取ることで、配列を渡すことができます。

int var = 100; // 型の宣言int *ptr = &var; // ポインタと参照の使用int **double_ptr = &ptr; // 二重ポインタの例printf(“ダブルポインタとポインタの逆参照：%d %d \n:”, **double_ptr, *ptr)int *ptr = 100; // int型のポインタの初期化

CUDA C

Colabへのリンク。

CUDAは、NVIDIAのGPUを利用して複雑な計算問題を並列化するためのコンピューティングプラットフォームです。Cの知識を活かしながら（冗談ですが）、CUDA Cコードの構造は、私たちが逐次進行した疑似コードの構造とまったく同じです。

CPU側では、アルゴリズムに何を期待するかを示すいくつかの定数をセットアップし、ライブラリをインポートし、変数を初期化し、いくつかのマクロを定義します。

#define NUM_ROWS 00000        // データセットのy次元、ブロックの数#define LINE_SIZE 100         // データセットのx次元#define NUM_ITERATIONS 100    // 最大イテレーション回数#define NUM_CLUSTERS 3        // k = 3 を実行しています#define MAX_INPUT_VALUE 100   // データの上限#define NUM_SEEDS 32          // シード/スレッドの数はLINE_SIZEの1/3です#define EPSILON_PERCENT 0.02  // 収束条件void initInputData(float** input) {    srand(1);     // データのメモリを割り当て... // mallocとrandを使用してデータを初期化    // GPU上でメモリを割り当て    cudaMalloc(input, NUM_ROWS * LINE_SIZE * sizeof(float));     // CPUのsample_dataからGPUのメモリにメモリをコピー    cudaMemcpy(*input, sample_data, NUM_ROWS * LINE_SIZE * sizeof(float), cudaMemcpyHostToDevice);    free(sample_data);}int main() {    float* inputData; // 入力データを初期化します。次元はNUM_ROWS x LINE SIZEです。    initInputData(&inputData); // デリファレンスして関数に渡す    // 出力のラベルとセントロイドを初期化します。    cudaExtent labelsExtent = make_cudaExtent(sizeof(int), LINE_SIZE, NUM_ROWS);    cudaPitchedPtr outputLabels; // 次の呼び出しに必要なポインタを作成します。    cudaMalloc3D(&outputLabels, labelsExtent); // GPU上にメモリを割り当てます。        cudaExtent centroidsExtent = make_cudaExtent(sizeof(float), NUM_CLUSTERS, NUM_ROWS);    cudaPitchedPtr outputCentroids; // 次の呼び出しに必要なポインタを作成します。    cudaMalloc3D(&outputCentroids, centroidsExtent); // GPU上にメモリを割り当てます。    cuda_kmeans <<<NUM_ROWS, NUM_SEEDS>>> (inputData, outputLabels, outputCentroids);    cudaDeviceSynchronize();        ... // デバイスからホストに出力をコピーします。}

違いを分解してみましょう。

主な関数は、ポインタを作成し、そのポインタのアドレスをinitInputDataに渡すことでデータの初期化を行います。関数は、ポインタをポインタ（float** input）へのポインタとして受け取ります。これにより、関数は元のポインタが保持するアドレスを変更することができます。入力は、cudaMallocを使用して初期化されたGPUメモリアドレスを指すようにポイントし、既にランダムな数値で満たされた一時的なホスト配列sample_dataからデータをコピーするcudaMemcpyを使用して埋められます。

次に、コードはk-means関数の結果を保持するためにデバイス上でメモリを割り当てます。関数はcudaExtentオブジェクトを作成するためにmake_cudaExtentを使用し、このオブジェクトの目的は多次元配列の次元をカプセル化することです。

cudaPitchedPointer型は、このピッチドメモリ空間にアドレス付けできるポインタを定義するために使用されます。このタイプのポインタは、cudaMalloc3Dによって割り当てられたメモリと共に使用するために特別に設計されており、cudaPitchedPtrおよびcudaExtentオブジェクトを受け取ってGPU上の線形メモリを割り当てます。

cuda_kmeans <<<NUM_ROWS, NUM_SEEDS>>> (inputData, outputLabels, outputCentroids);

GPUコードに入ると、各ブロックがデータの行に対応し、各スレッドがシードに対応するようにグリッドを定義します。

シードは各ブロックの単一のスレッドによって初期化され、完全に異なるシードが確保されます。

__global__ void cuda_kmeans(float* input, cudaPitchedPtr outputLabels, cudaPitchedPtr outputCentroids) {    int row = blockIdx.x;    int seed = threadIdx.x;    // shared memory is shared between threads in blocks    __shared__ float input_shm[LINE_SIZE];    __shared__ float error_shm[NUM_SEEDS];    __shared__ float seed_centroids[NUM_SEEDS][NUM_CLUSTERS];    __shared__ int seed_labels[NUM_SEEDS][LINE_SIZE];        ... // get a single row of data    ... // populate input_shm    ... // populating the struct core_params    // the actual k-means function    kmeans(core_params);     // find seed with smallest error    calcError(core_params);    __syncthreads();    if (seed == 0) {        int* labels_line = LABELS_LINE(outputLabels, row);        float* centroids_line = CENTROIDS_LINE(outputCentroids, row);        labels_line[threadIdx.x] = seed_labels[seed][threadIdx.x];        centroids_line[threadIdx.x] = seed_centroids[seed][threadIdx.x];    }}

CUDA Cコードでは、データ、セントロイド、ラベル、およびエラーに対して共有メモリが使用されます。ただし、Pythonとは異なり、コードは共有メモリへのポインタを取り、それらを構造体に格納します。これは単に変数をまとめて渡すための方法です。最後に、cuda_kmeansは実際のk-meansアルゴリズムを呼び出し、core_paramsを渡します。

__device__ void kmeans(core_params_t& core_params) {    DECLARE_CORE_PARAMS(core_params);    getInitialCentroids(core_params);    sort_centroids(centroids, num_clusters);    float yard_stick = findYardStick(core_params);    float* oldCentroids = (float*)malloc(NUM_CLUSTERS * sizeof(float));    struct work_params_t work_params;    work_params.min = find_min(line, LINE_SIZE);    work_params.max = find_max(line, LINE_SIZE);    work_params.aux_buf1 = (int*)malloc(NUM_CLUSTERS * sizeof(int));    work_params.aux_buf2 = (int*)malloc(NUM_CLUSTERS * sizeof(int));    work_params.aux_buf3 = (float*)malloc(NUM_CLUSTERS * sizeof(float));    for (int iterations = 0; true; iterations++) {        bool stop = (iterations > 100) || (iterations > 0 && (converged(core_params, oldCentroids, yard_stick * EPSILON_PERCENT)));        if (stop)            break;        memcpy(oldCentroids, core_params.centroids, NUM_CLUSTERS * sizeof(float));        getLabels(core_params);        getCentroids(core_params, work_params);    }    free(work_params.aux_buf1);    free(work_params.aux_buf2);    free(work_params.aux_buf3);    free(oldCentroids);}

デバイス関数では、まずcore_params構造体から変数に値を取り出すためにDECLARE_CORE_PARAMSマクロを使用します。

次に、Pythonと同じk-meansアルゴリズムを実行しますが、変数の代わりに構造体を渡し、メモリ、ポインタ、および型を管理する必要があるという違いがあります。

ベンチマーク

非並列化されたk-meansと比較するために、scikit-learnのk-meansモジュールをインポートします。

ベンチマークでは、3つのクラスタを持つ100,000行×100列のデータを処理します。scikit-learnには異なる行に対して並列処理のk-meansがないので、forループ内で行を順次実行します。

Colabでのベンチマークでは、無料のT4 GPU Colabインスタンスを使用します。

結果は良好です – PythonのNumbaコードは非並列化されたCPUコードよりも100倍高速であり、CUDA Cコードは1000倍高速です。カーネル関数は簡単にスケーラブルであり、アルゴリズムは高次元のクラスタリングをサポートするために変更できます。

なお、CおよびPythonの両方のアルゴリズムでのランダムな初期重心の生成は、すべてのコアを使用するために最適化されていません。改善された場合、PythonのアルゴリズムはCのコードの実行時間に近づくかもしれません。

異なるデータセットでk-means関数を100回実行し、結果の時間を記録すると、最初のイテレーションはCとPythonのコンパイルにかかる時間のために著しく遅くなることに気付きます。

結論

これで独自のカスタムGPUカーネルを書く準備が整いました！残る質問は、並列化されたデータ処理のワークロードでCUDA CまたはNumbaを使用するべきかどうかです。それは状況によります。どちらも市販のscikit-learnよりもはるかに高速です。私の場合、CUDA Cのバッチ処理されたk-meansの実装は、Numbaを使用した同等の実装よりも約3.5倍高速でしたが、Pythonは可読性やPythonで主に作業するチームにおける特殊なCプログラミングスキルへの依存の軽減など、重要な利点を提供しています。さらに、特定の実装の実行時間は、GPU上でシリアル化された操作をトリガーしないようにコードを最適化するなど、どれくらい効率化されているかにも依存します。結論として、Cと並列プログラミングの両方が初めてである場合、アルゴリズムのプロトタイプを作成するためにNumbaを使用し、追加の高速化が必要な場合はCUDA Cに変換することをお勧めします。

参考文献

1. Scikit-learn：Pythonによる機械学習、Pedregosa et al.、JMLR 12、pp. 2825-2830、2011。
2. NVIDIA、Vingelmann、P.＆Fitzek、F.H.P.、2020. CUDA、リリース：10.2. 89、入手先：https://developer.nvidia.com/cuda-toolkit.
3. Lam、Siu Kwan、Antoine Pitrou、およびStanley Seibert。“Numba: A llvm-based python jit compiler.”第2回LLVM Compiler Infrastructure in HPCワークショップの論文。2015年。
4. Harris、C.R.、Millman、K.J.、van der Walt、S.J. et al. “Array programming with NumPy.” Nature 585, 357–362 (2020). DOI：10.1038/s41586–020–2649–2。（出版者リンク）。

We will continue to update VoAGI; if you have any questions or suggestions, please contact us!