OpenCLを用いたMQL5におけるCPUからGPUへの実践的移行パス

はじめに

MQL5におけるCPUからGPUへの移行は、一見すると単純なステップのように見えます。つまり、GPUの方が高速に計算できるのであれば、取引分析も自動的に高速化されるはずだという考え方です。しかし実際には、問題はもっと複雑です。GPUは確かに大きな性能向上をもたらす可能性がありますが、それはタスクが並列計算モデルに適している場合に限られます。そうでない場合、得られるのは速度向上ではなく、複雑化したアーキテクチャと同等もしくはそれ以上のコストだけになる可能性があります。

これはアルゴリズム取引において特に重要です。市場データの分析、パラメータの総当たり、仮説の大規模検証、繰り返しパターンの探索などは、大量の計算を必要とします。こうした場面こそ、GPUの真価が発揮される領域です。GPUは、同じ操作を多数の要素に対して並列に実行でき、結果を後からまとめて取得できるような処理に強みを持っています。このようなケースでは、グラフィックスカードは単なる飾りではなくなり、完全に独立した本格的な計算リソースとしての役割を持つようになります。

しかし、GPUにはコストも伴います。計算前にはデータ準備、デバイスへの転送、カーネル実行待機、結果の回収が必要です。小規模な処理では、このオーバーヘッドが重くなり、CPUの方が効率的である場合もあります。特にデータ量が少ない場合やロジックが頻繁に変化する場合には、GPU移行はむしろ非効率になることもあります。場合によっては、特にタスクが頻繁に変更される場合、柔軟な条件分岐が必要な場合、または少量のデータに関連する場合、それが邪魔になることもあります。

MQL5環境では、OpenCLがアプリケーションロジックとGPUをつなぐ役割を担います。これにより、計算負荷の高い部分をメインプログラムの外に移し、GPU上でバッチデータ処理を構成することが可能になります。しかしOpenCL自体は高速化の魔法ではなく、最初から並列計算を前提としたアーキテクチャ設計と、CPU-GPU間のデータ交換の最小化があって初めて効果を発揮します。

本記事ではGPUを、計算パイプライン内の独立した処理層として扱います。これは、負荷が大きく反復的な処理を対象としています。このアプローチは、研究、最適化、パターン探索といったタスクにおいて有効です。これらの領域では、計算量の増加速度が、人間が許容できる待ち時間を上回ります。ここでの実践的なポイントは単純です。まず「何をGPUに移す価値があるのか」を正確に理解し、その上で初めて実際の高速化効果を期待すべきだということです。

環境の準備

OpenCLの作業は、計算ではなく準備から始まります。まずプログラムは利用可能なデバイスを見つけ、ワーキングコンテキストを作成し、カーネルを準備し、データ用メモリを確保します。これらは一見すると単なる技術的な形式作業のように見えますが、この段階で大きな生産性の損失が発生することがよくあります。

最も大きな誤りは、GPUを通常の関数呼び出しのように扱うことです。呼び出して結果を受け取り、そのまま次に進むという考え方です。しかし実際には、その呼び出しの裏側には一連の処理があります。コンテキストの作成、プログラムの準備、カーネルのコンパイル、メモリ割り当て、データ転送、そしてようやく計算開始という流れです。これを何度も繰り返すと、GPUは計算よりも準備に多くの時間を費やすことになります。

そのため、良い実装では可能な限り多くの処理を一度だけ実行します。コンテキストは事前に作成され、その後は再利用されます。カーネルもコードが変わらない限り一度だけコンパイルされます。メモリバッファも必要がない限り再生成せず、再利用することが望ましいです。このアプローチはオーバーヘッドを削減し、GPUの利用を実用的なものにします。

データ転送にも同じことが当てはまります。GPUは、頻繁に小さなタスクを往復させる用途には適していません。この方法では計算時間ではなくデータ交換に時間が費やされます。より効率的なのは、データを大きなバッチとしてまとめ、計算回数を減らしつつ一回あたりの負荷を高くすることです。一度の大きな実行は、ほとんどの場合、小さな実行の繰り返しよりも優れています。

もう一つの一般的な問題は、不要な同期です。各ステップごとにプログラムが停止してGPUの完了を待つ場合、デバイスはアイドル状態になります。これは全体の加速効果を低下させます。GPUには可能な限り連続的に作業を行わせ、結果は必要なときにのみ取得するように設計する方がよいです。

これはMQL5において特に重要です。取引プログラムは遅延に敏感であり、設計上の雑さはすぐに影響として現れます。コンテキストが常に再生成され、メモリが無秩序に確保され、カーネルが毎回コンパイルされるような設計では、GPUは加速装置ではなく遅延の原因になります。

したがって、ここでの基本原則は非常に単純です。事前に準備できるものはすべて事前に準備すること。再利用可能なものは再生成しないこと。この規律が保たれている場合、OpenCLは重い計算の高速化に確かに役立ちます。この規律が欠けている場合、その利点は計算が始まる前に失われてしまいます。

メモリの処理

GPUというと、多くの人はまず計算性能を思い浮かべます。しかし実際には、計算速度そのものよりも「データをどのように供給するか」が結果を左右することが多いです。どれほど高速なGPUであっても、メモリの扱いが悪ければ良い結果は得られません。

OpenCLにおいてメモリは、単なるデータ保存領域ではありません。性能はその使い方に直接依存します。GPUには複数のメモリ階層があり、高速に動作する領域と低速な領域が存在します。小さな内部メモリは高速にアクセスできますが、デバイスのメインメモリへのアクセスははるかにコストが高くなります。

これにより重要な経験則が導かれます。GPUが低速メモリへアクセスする回数をできるだけ減らすようにデータを設計すべきです。不要な読み書きが少ないほど性能は向上します。特に重要なのは、CPUとGPU間で同じデータを何度も不必要に転送しないことです。多くの場合、ボトルネックになるのは計算ではなくデータ転送です。

簡単に言えば、GPUは「大量のデータ配列を一度送信し、同様の処理を多数実行し、その後結果をまとめて返す」ようなタスクを好みます。一方で、「小さなデータを頻繁に送り、そのたびに応答を待つ」ような処理は好みません。このようなモードでは、加速効果はすぐに失われます。

これは特に取引用途において重要です。プログラムが毎ステップごとにデータを準備し、転送し、結果を待ち、また同じ処理を繰り返すような構造では、性能は不安定になります。より良い方法は、必要なデータを事前にまとめて準備し、単一のブロックとしてデバイスに送信し、計算を実行し、その結果をメインプログラム側で利用することです。

ここでは役割分担が重要です。CPUは制御センターとして機能し、データを収集し、計算を開始し、最終判断を行います。一方GPUは、多数の同一処理を高速に繰り返す部分を担当します。この構造は、すべてを単一デバイスに任せるよりも、ほとんどの場合でより安定かつ効率的です。

また、典型的な誤りとして「カーネル起動回数が多すぎる」問題があります。一見すると、新しいデータが来るたびにGPUへ送るのは合理的に見えます。しかし、カーネル起動自体にもコストが発生します。起動回数が多すぎると、そのオーバーヘッドが性能向上を打ち消してしまいます。そのため多くの場合、GPUの呼び出し回数は少なくし、その代わり一度に与えるタスクを大きく・単純にする方が適切です。

結果として、メモリ管理は単なる技術的な細部ではなく、パフォーマンスを決定する主要因の一つになります。データが適切に整理されていれば、GPUは効率よく処理を流せるようになります。そうでなければ、どれほど高性能なデバイスであっても待機や不要な転送に時間を浪費することになります。

結論：GPUは、データが正しく供給されて初めて計算を高速化します。不必要な転送、呼び出し、小さな起動が少ないほど、実際の性能効果は高くなります。

プログラムの構築

アプリケーションプログラムにおいて、CPUとGPUは競合するのではなく、協調して動作します。CPUは制御センターとしての役割を維持し、初期データを生成し、必要なメソッドを呼び出し、結果を受け取り、実行時間を比較します。一方でGPUは計算処理のみを担当します。これは典型的な構造です。プログラム全体をデバイスへ移すのではなく、実際に並列性が存在する部分だけをGPUに与えるべきです。

OpenCLを用いたプログラム構築を理解する最も簡単な方法は、具体例を使用することです。MQL5の標準ライブラリには、行列乗算の実装例があります。メインプログラムでは、まず2つの行列が作成され、ランダムな値で初期化されます。

void OnStart()
  {
//--- matrix A 1000x2000
   int rows_a = 1000;
   int cols_a = 2000;
//--- matrix B 2000x1000
   int rows_b = cols_a;
   int cols_b = 1000;
//--- matrix C 1000x1000
   int rows_c = rows_a;
   int cols_c = cols_b;
//--- matrix A: size=rows_a*cols_a
   int size_a = rows_a * cols_a;
   int size_b = rows_b * cols_b;
   int size_c = rows_c * cols_c;
//--- prepare matrix A
   float matrix_a[];
   ArrayResize(matrix_a, rows_a * cols_a);
   for(int i = 0; i < rows_a; i++)
      for(int j = 0; j < cols_a; j++)
        {
         matrix_a[i * cols_a + j] = (float)(10 * MathRand() / 32767);
        }
//--- prepare matrix B
   float matrix_b[];
   ArrayResize(matrix_b, rows_b * cols_b);
   for(int i = 0; i < rows_b; i++)
      for(int j = 0; j < cols_b; j++)
        {
         matrix_b[i * cols_b + j] = (float)(10 * MathRand() / 32767);
        }

まずCPU上で逐次計算が実行され、その後、同じ計算がGPU上でも実行されます。

//--- CPU: calculate matrix product matrix_a*matrix_b
   float matrix_c_cpu[];
   ulong time_cpu = 0;
   if(!MatrixMult_CPU(matrix_a, matrix_b, matrix_c_cpu, rows_a, cols_a, cols_b, time_cpu))
     {
      PrintFormat("Error in calculation on CPU. Error code=%d", GetLastError());
      return;
     }
//--- calculate matrix product using GPU
   float matrix_c_gpu_method1[];
   float matrix_c_gpu_method2[];
   ulong time_gpu_method1 = 0;
   ulong time_gpu_method2 = 0;
   if(!MatrixMult_GPU(matrix_a, matrix_b, matrix_c_gpu_method1, matrix_c_gpu_method2,
       rows_a, cols_a, cols_b, size_a, size_b, size_c, time_gpu_method1, time_gpu_method2))
     {
      PrintFormat("Error in calculation on GPU. Error code=%d", GetLastError());
      return;
     }

計算そのものは、それぞれ独立したメソッドで実行されます。GPU側の処理は、同一問題に対する2つの実装として提示されており、それはナイーブ実装と最適化実装です。これにより、違いを言葉ではなく、コードおよび実行時間の観点から直接確認することができます。

まずCPU版を見てみましょう。ここではすべてが非常に明確で、典型的な三重ループによる実装です。結果行列の各要素は逐次的に計算されます。

bool MatrixMult_CPU(const float &matrix_a[], const float &matrix_b[], float &matrix_c[],
                    const int rows_a, const int cols_a, const int cols_b, ulong &time_cpu)
  {
   int size = rows_a * cols_b;
   if(ArrayResize(matrix_c, size) != size)
      return(false);
//--- CPU calculation started
   time_cpu = GetMicrosecondCount();
   for(int i = 0; i < rows_a; i++)
     {
      for(int j = 0; j < cols_b; j++)
        {
         float sum = 0.0;
         for(int k = 0; k < cols_a; k++)
           {
            sum += matrix_a[cols_a * i + k] * matrix_b[cols_b * k + j];
           }
         matrix_c[cols_b * i + j] = sum;
        }
     }
//--- CPU calculation finished
   time_cpu = ulong((GetMicrosecondCount() - time_cpu) / 1000);
//---
   return(true);
  }

このコードは、タスクの基本的な考え方を示しています。2つの行列があり、行と列の積の総和が計算されます。処理は逐次的に実行されます。CPU上ではこれは透過的に動作し、余分な準備も必要ありません。しかし、まさにここに主要な制約があります。行列サイズが大きくなるにつれて、逐次計算は急速にコストの高い処理になります。

次にGPU側の処理に移ります。ここで重要なのは基本原則です。この例ではOpenCLの処理は別メソッドに隠蔽されており、メインプログラムはクリーンな状態を保ち、GPUに関するすべての処理は1か所に集約されています。

bool MatrixMult_GPU(const float &matrix_a[], const float &matrix_b[], float &matrix1_c[], float &matrix2_c[],
                    const int rows_a, const int cols_a, const int cols_b, const int size_a, const int size_b,
                    const int size_c, ulong &time1_gpu, ulong &time2_gpu)
  {
   const int task_dimension = 2;
//--- prepare matrices for result
   if(ArrayResize(matrix1_c, size_c) != size_c || ArrayResize(matrix2_c, size_c) != size_c)
      return(false);
   ArrayFill(matrix1_c, 0, size_c, (float)0.0);
   ArrayFill(matrix2_c, 0, size_c, (float)0.0);

ここでは重要な点がすぐに分かります。2つのGPU実装用の結果行列は事前に準備されています。これにより、計算処理とメモリ準備が混在することを防いでいます。この段階では、プログラムはすでに規律ある構造になっており、まず配列の確保がおこなわれ、その後にOpenCLの処理が続きます。

次に、OpenCLコンテキストが作成され、初期化されます。

//--- OpenCL
   ulong timei_gpu = GetMicrosecondCount();
   COpenCL OpenCL;
   if(!OpenCL.Initialize(cl_program, true))
     {
      PrintFormat("Error in OpenCL initialization. Error code=%d", GetLastError());
      return(false);
     }

ここでOpenCLの実践的な意味が明確になります。ここまでの段階では、プログラムは通常のMQL5コードでした。しかしここからは、GPU向けの計算環境を準備しています。そして特に重要なのは、これは無料の処理ではないという点です。初期化時間は別途測定されます。高速化について議論する前に、まず環境の起動そのものにどれだけコストがかかるのかを正確に把握する必要があります。

次に、2つのカーネルが作成されます。1つ目は単純な並列化バージョンです。2つ目はローカルグループを用いたより高度な実装です。

//--- create kernels
   OpenCL.SetKernelsCount(2);
   OpenCL.KernelCreate(0, "MatrixMult_GPU1");
   OpenCL.KernelCreate(1, "MatrixMult_GPU2");

ここでは、実行時間が主にOpenCLプログラム内で使用されるアルゴリズムの品質に依存するという重要な点に注意する必要があります。単純に計算を並列化することもできますし、メモリ管理を改善することもできます。この例ではその両方のアプローチが含まれており、それによって違いがより明確に示されています。

次にバッファを準備します。入力となる行列はデバイスへコピーされ、結果用のバッファは別途作成されます。

//--- create buffers
   OpenCL.SetBuffersCount(3);
//---
   if(!OpenCL.BufferFromArray(0, matrix_a, 0, size_a, CL_MEM_READ_ONLY))
     {
      PrintFormat("Error in BufferFromArray for matrix A. Error code=%d", GetLastError());
      return(false);
     }
   if(!OpenCL.BufferFromArray(1, matrix_b, 0, size_b, CL_MEM_READ_ONLY))
     {
      PrintFormat("Error in BufferFromArray for matrix B. Error code=%d", GetLastError());
      return(false);
     }
   if(!OpenCL.BufferCreate(2, size_c * sizeof(float), CL_MEM_WRITE_ONLY))
     {
      PrintFormat("Error in BufferCreate for matrix C. Error code=%d", GetLastError());
      return(false);
     }

これはすでに実際に動作するGPU構造になっています。データはデバイスへ転送され、そこで計算が行われ、その後ホスト側へ戻されます。この順序こそが高速化を可能にするものです。これを細かく分割してしまうと、GPUは計算そのものではなく、準備やデータ交換に多くの時間を費やすようになってしまいます。

その後、最初のカーネルの引数が設定されます。

//--- prepare arguments for kernel 0
   int kernel_index = 0;
   OpenCL.SetArgumentBuffer(kernel_index, 0, 0);
   OpenCL.SetArgumentBuffer(kernel_index, 1, 1);
   OpenCL.SetArgumentBuffer(kernel_index, 2, 2);
   OpenCL.SetArgument(kernel_index, 3, rows_a);
   OpenCL.SetArgument(kernel_index, 4, cols_a);
   OpenCL.SetArgument(kernel_index, 5, cols_b);
   timei_gpu = ulong((GetMicrosecondCount() - timei_gpu) / 1000);
   PrintFormat("time of initialization GPU =%d ms", timei_gpu);

ここでのロジックは非常にシンプルです。カーネルはデータバッファと行列サイズを受け取ります。OpenCLカーネル内部では、どのスレッドがどの要素を担当するかが決定されます。ここは重要なポイントです。GPUはそれ自体では配列の意味や構造を解釈することはできません。そのため、この構造は明示的に伝える必要があります。

その後、問題サイズが設定され、最初の計算オプションが実行されます。

//--- set task dimension a_rows x b_cols
   uint global_work_size[2];
//--- set dimensions
   global_work_size[0] = rows_a;
   global_work_size[1] = cols_b;
   uint global_work_offset[2] = {0, 0};
//--- GPU calculation start kernel 0
   time1_gpu = GetMicrosecondCount();
   if(!OpenCL.Execute(kernel_index, task_dimension, global_work_offset, global_work_size))
     {
      PrintFormat("Error in Execute. Error code=%d", GetLastError());
      return(false);
     }
   if(!OpenCL.BufferRead(2, matrix1_c, 0, 0, size_c))
     {
      PrintFormat("Error in BufferRead for matrix1 C. Error code=%d", GetLastError());
      return(false);
     }
//--- GPU calculation finished
   time1_gpu = ulong((GetMicrosecondCount() - time1_gpu) / 1000);

これが最初のGPUオプションです。これは基本的な考え方を示しており、1つのスレッドが1つの出力要素を計算します。この構造だけでも並列性は得られますが、まだメモリ最適化の可能性を十分には活用していません。そのため、この例ではもう1つのカーネルが用意されています。こちらも同様に引数が設定されますが、実行方法が異なり、ローカルワークグループを使用します。

//--- prepare arguments for kernel 1
   kernel_index = 1;
//--- set arguments
   OpenCL.SetArgumentBuffer(kernel_index, 0, 0);
   OpenCL.SetArgumentBuffer(kernel_index, 1, 1);
   OpenCL.SetArgumentBuffer(kernel_index, 2, 2);
   OpenCL.SetArgument(kernel_index, 3, rows_a);
   OpenCL.SetArgument(kernel_index, 4, cols_a);
   OpenCL.SetArgument(kernel_index, 5, cols_b);
   uint local_work_size[2];
   local_work_size[0] = BLOCK_SIZE;
   local_work_size[1] = BLOCK_SIZE;

ここで実践的な違いが現れます。最初のバージョンは単純に計算を並列化するだけです。一方、2つ目のバージョンでは計算をブロック単位で整理しています。これは一見した以上に重要な違いです。GPUは並列性だけでなく、効率的なメモリ構成も重視します。そのため、ブロック化とローカルワークグループは大きな効果をもたらします。

2つ目のカーネル実行では、ワークグループの次元も追加で指定されます。

//--- GPU calculation start, kernel1
   time2_gpu = GetMicrosecondCount();
   if(!OpenCL.Execute(kernel_index, task_dimension, global_work_offset, global_work_size, local_work_size))
     {
      PrintFormat("Error in Execute. Error code=%d", GetLastError());
      return(false);
     }
   if(!OpenCL.BufferRead(2, matrix2_c, 0, 0, size_c))
     {
      PrintFormat("Error in BufferRead for matrix2 C. Error code=%d", GetLastError());
      return(false);
     }
//--- GPU calculation finished
   time2_gpu = ulong((GetMicrosecondCount() - time2_gpu) / 1000);
//--- remove OpenCL objects
   OpenCL.Shutdown();
//---
   return(true);
  }

ここからが最も興味深い部分です。実際にOpenCLコード内部で何が行われているのかを見ていきます。最初のカーネル実装は、可能な限りシンプルな構造になっています。

__kernel void MatrixMult_GPU1(__global float *matrix_a,
                              __global float *matrix_b,
                              __global float *matrix_c,
                              int rows_a, int cols_a, int cols_b)
  {
   int i = get_global_id(0);
   int j = get_global_id(1);
   float sum = 0.0;
   for(int k = 0; k < cols_a; k++)
     {
      sum += matrix_a[cols_a * i + k] * matrix_b[cols_b * k + j];
     }
   matrix_c[cols_b * i + j] = sum;
  }

これは、数式上のロジックをGPUへほぼそのまま移したような実装です。各スレッドは自身の座標を受け取り、結果行列の1要素を計算します。

2つ目のバージョンは、さらに興味深いものになっています。こちらではローカル配列とスレッド同期が使用されています。

__kernel void MatrixMult_GPU2(__global float *matrix_a,
                              __global float *matrix_b,
                              __global float *matrix_c,
                              int rows_a, int cols_a, int cols_b)
  {
   int group_i = get_group_id(0);
   int group_j = get_group_id(1);
   int i = get_local_id(0);
   int j = get_local_id(1);
   __local float submatrix_a[BLOCK_SIZE][BLOCK_SIZE];
   __local float submatrix_b[BLOCK_SIZE][BLOCK_SIZE];
   int offset_b = BLOCK_SIZE * group_i;
   int offset_a_start = cols_a * BLOCK_SIZE * group_j;
   float sum = (float)0.0;

ここでは、計算構造がすでに異なっていることが分かります。スレッドはグループ化され、データはブロック内部のメモリへ読み込まれます。これによりグローバルメモリへのアクセス回数が減少し、デバイスをより効率的に動作させることができます。

次に、データ断片が読み込まれ、スレッドの同期が行われます。

   for(int offset_a = offset_a_start;
       offset_a < offset_a_start + cols_a;
       offset_a += BLOCK_SIZE,
       offset_b += BLOCK_SIZE * cols_b)
     {
      submatrix_a[i][j] = matrix_a[offset_a + cols_a * i + j];
      submatrix_b[i][j] = matrix_b[offset_b + cols_b * i + j];
      barrier(CLK_LOCAL_MEM_FENCE);
      for(int k = 0; k < BLOCK_SIZE; k++)
         sum += submatrix_a[i][k] * submatrix_b[k][j];
      barrier(CLK_LOCAL_MEM_FENCE);
     }

GPU実装が2種類存在する理由に対する回答はこちらです。1つ目は並列化を示しています。2つ目はメモリ最適化を示しています。そして実際の開発現場では、まさにこれが高速化の成否を左右することがよくあります。

最終的な出力結果は、出力配列へ書き戻されます。

   int offset_c = BLOCK_SIZE * (cols_b * group_j + group_i);
   matrix_c[offset_c + cols_b * i + j] = sum;
  };

これら2つのカーネルは、メインプログラム内で実行時間を比較され、さらにCPU実装との比較によって計算精度の検証もおこなわれます。

//--- calculate CPU/GPU ratio
   double CPU_GPU_ratio1 = 0;
   double CPU_GPU_ratio2 = 0;
   if(time_gpu_method1 != 0)
      CPU_GPU_ratio1 = 1.0 * time_cpu / time_gpu_method1;
   if(time_gpu_method2 != 0)
      CPU_GPU_ratio2 = 1.0 * time_cpu / time_gpu_method2;
   PrintFormat("time CPU=%d ms, time GPU global work groups =%d ms, CPU/GPU ratio: %f",
                                           time_cpu, time_gpu_method1, CPU_GPU_ratio1);
   PrintFormat("time CPU=%d ms, time GPU local work groups  =%d ms, CPU/GPU ratio: %f",
                                           time_cpu, time_gpu_method2, CPU_GPU_ratio2);
   PrintFormat("time matrix CPU=%d ms", time_mat);

補足として、組み込みの行列演算の使用も比較対象として追加しておきましょう。

//--- matrix
   matrix<float> A, B, C;
   if(!A.Assign(matrix_a) || !B.Assign(matrix_b))
     {
      PrintFormat("Error of copy data to matrices. Error code=%d", GetLastError());
      return;
     }
   if(!A.Reshape(rows_a, cols_a) || !B.Reshape(rows_b, cols_b))
     {
      PrintFormat("Error of copy data to matrices. Error code=%d", GetLastError());
      return;
     }
   ulong time_mat = GetMicrosecondCount();
   C = A.MatMul(B);
   time_mat = ulong((GetMicrosecondCount() - time_mat) / 1000); 

実験では、ハイブリッドシステム上における行列乗算時間が比較されました。対象はCPUと、2種類のGPU（NVIDIA GeForce RTX 4060 Laptop GPUとIntel Iris Xe Graphics）です。ベースラインとしては単純なCPU実装が使用され、その実行時間はおよそ2056～2180ミリ秒でした。ここで注目すべきなのは、最適化された行列演算であり、こちらは約32～34ミリ秒という安定した実行時間を示しました。このクラスのCPUとしては、ベクトル化処理による期待性能に近い結果であると考えられます。

計算をGPUへ移行した場合、結果は一様ではなく、OpenCLの効率が計算ブロックの構成にどれほど強く依存しているかが明確に示されました。RTX 4060では、グローバルワークグループを使用した場合の実行時間は約38ミリ秒であり、CPUによる最適化済み行列演算に対して明確な優位性は見られませんでした。しかし、ローカルワークグループへ切り替えることで状況は劇的に変化し、実行時間は11ミリ秒まで短縮されました。これは、ローカルメモリ内でデータを再利用することでグローバルメモリへの負荷を軽減し、GPU計算ユニットをより効率的に利用するタイル化アプローチが十分に機能していることを示しています。

同様の傾向は、統合GPUであるIntel Iris Xeでも観測されましたが、その差はさらに顕著でした。グローバルグループモードでは実行時間は213ミリ秒でしたが、ローカルグループを使用すると45ミリ秒まで短縮されました。離散GPUに対する全体的な性能差は依然として存在するものの、相対的な高速化率はこちらの方がより顕著であり、性能の低いGPUほどメモリアクセス最適化への感度が高いことを示しています。

GPU初期化時間にも特に注意を払う必要があります。RTX 4060では約99ミリ秒であったのに対し、Iris Xeでは約4ミリ秒でした。この要素はアプリケーションシナリオにおいて無視できず、計算パイプライン全体の効率に直接影響します。

全体として、これらの結果は、メモリ帯域制約型の実行から計算効率重視の実行モードへの移行を示す典型的な例となっています。この規模のタスクではCPUも依然として競争力を維持していますが、GPUはローカル計算ユニットを適切に構成した場合にのみ、その優位性を発揮します。これは特にRTX 4060で顕著であり、非最適化実装と最適化実装の差は約3～4倍にも達しました。この差こそが、グラフィックスアクセラレータを非効率に利用している状態と、十分活用している状態との境界を事実上定義していると言えます。

ローソク足パターンのテスト

行列乗算の例は、GPUがどのように計算を高速化できるかをよく示しています。しかし、トレーダーにとってはそれだけでは不十分です。この手法が、実際に市場に関連する問題に適用できるかどうかを理解することの方が重要です。そこで次のステップとして、抽象的な数学からローソク足パターン解析へと移行します。

ローソク足パターンは通常、よく知られた名前を持つ完成された形として説明されます。これらは図としては認識しやすく、書籍の中でも説得力のある形で提示されることが多いです。しかし、より厳密に考えると、こうしたモデルが実際に統計的にどの程度裏付けられているのかという疑問が生じます。正確な基準はどこにあるのでしょうか。あるパターンが、理論上ではなく実データ上で有効であることをどのように判断するのでしょうか。

あらかじめ教科書的な既知の図形を探すのではなく、別のアプローチを取ることができます。あらかじめ用意されたテンプレートを市場に当てはめるのではなく、過去に同様の状況でどのように振る舞ってきたかを調べるのです。

直近の数本のローソク足を取り出します。これが現在の市場状況、すなわち参照パターンとなります。その後、プログラムは過去データを走査し、その形状が現在のパターンと類似している領域を探します。比較はパターン名ではなく、実際のローソク足の特徴、つまり実体、上ヒゲ、下ヒゲに基づいておこなわれます。さらに、市場はほとんど同じ形を完全には繰り返さないため、わずかな誤差も許容されます。

過去に類似した区間が見つかった場合、プログラムは推測をおこなうのではなく、その結果を検証します。その状況に対して、テイクプロフィットとストップロスのレベル、さらに時間制限を設定して取引をシミュレーションします。そして、そのようなケースが過去において利益で終わったのか損失で終わったのかを計算します。

ここでOpenCLが特に有用になります。このタスクは多数の類似したチェックから構成されています。つまり、過去の多くの区間を走査し、それらを現在のテンプレートと比較し、それぞれについて取引結果を計算する必要があります。CPUでも実行は可能ですが、大量のデータでは計算が重くなります。一方GPUにとっては自然な負荷であり、多数の独立した計算を並列に処理できます。

さらに、この設計では単一の取引パラメータの組み合わせだけでなく、パラメータのグリッド全体を同時にテストします。つまり、テイクプロフィットとストップロスのさまざまな値を同時に評価します。これにより、どのパラメータが同様の市場条件で最も妥当であったかを評価できます。

まずはOpenCL側のロジックから始めます。この段階でタスクは実際の形を持ち始めます。CPUはこの設計において指揮役として残りますが、重い処理（検索、比較、シミュレーション）はすべてGPUに任されます。

ここにあるのは分析パイプラインです。

__kernel void PatternStats3D(__global const float4 *price,
                             __global const float  *tp,
                             __global const float  *sl,
                             __global float        *global_stats,
                             const int bars,
                             const float tolerance,
                             const int horizon)
  {
   const int lid = get_local_id(0);
   const int itp = get_global_id(1);
   const int isl = get_global_id(2);
   const int total_loc = get_local_size(0);
   const int tp_count = get_global_size(1);
   const int sl_count = get_global_size(2);

ソースデータは極めてコンパクトな形で構成されています。履歴データはfloat4ベクトル配列として渡されます。各エントリは1本のローソク足であり、open、high、low、closeを含みます。これは重要です。データを別々の配列に分割せず、アクセス構造を複雑化していません。GPUは密なデータ構造を扱う方が効率的であり、この点が最大限活用されています。

tpとslの配列は別々に渡されます。このようにすることで、タスクの第2次元と第3次元を即座に定義しています。各スレッドは履歴の一部分だけでなく、特定のトレードパラメータの組み合わせとも対応します。その結果、計算空間は三次元構造（history × TP × SL）になります。

その後、実際の処理が開始されます。各ワークグループ内のスレッドはそれぞれ固有のlidを取得し、グループサイズに等しいステップで履歴データを走査していきます。

   __local int buf_stat[BLOCK_SIZE][STAT_DIM];
   int local_stat[STAT_DIM];
   for(int i=0;i<STAT_DIM;i++)
     local_stat[i]=0;
//---
   float4 pattern[PATTERN_SIZE];
   if(bars < (PATTERN_SIZE + horizon))
      return;
   for(int i = 0; i < PATTERN_SIZE; i++)
      pattern[i] = price[bars - 1 - PATTERN_SIZE + i];
//--- border
   for(int i = lid; i < (bars - horizon - PATTERN_SIZE); i += total_loc)
     {
      bool match = true;
      //--- pattern check
      for(int k = 0; k < PATTERN_SIZE ; k++)
        {
         float4 a = price[i + k];
         float body_a  = a.w - a.x;
         float body_b  = pattern[k].w - pattern[k].x;
         if(fabs(body_a - body_b) > tolerance)
           {
            match = false;
            break;
           }
         float upper_a = a.y - fmax(a.w, a.x);
         float upper_b = pattern[k].y - fmax(pattern[k].w, pattern[k].x);
         if(fabs(upper_a - upper_b) > tolerance)
           {
            match = false;
            break;
           }
         float lower_a = fmin(a.w, a.x) - a.z;
         float lower_b = fmin(pattern[k].w, pattern[k].x) - pattern[k].z;
         if(fabs(lower_a - lower_b) > tolerance)
           {
            match = false;
            break;
           }
        }

これは典型的な手法です。各ローソク足ごとにフローを作成することはしません。それはコストが高すぎるためです。代わりに、各スレッドは自身のデータストリップを処理します。これにより、余計なカーネル起動なしで、負荷が均等に分散されます。

処理開始前に、基準となるパターンを構築します。履歴データから直近のローソク足を取り出します。これが現在の市場状態であり、対象となる参照情報です。外部のパターンは使いません。推測もおこないません。実際の価格状態のみを扱います。

次に重要なポイントである比較処理に移ります。履歴上の各位置について、その区間が基準と類似しているかどうかを確認します。比較はローソク足の構造に基づいておこなわれます。

• 実体
• 上ヒゲ
• 下ヒゲ

そしてこれらはすべて許容誤差付きで評価されます。この点は重要なニュアンスです。完全一致は要求しません。市場は同一形状を繰り返すわけではないため、重要なのはピクセル単位の一致ではなく「形状」です。

いずれかの要素が許容範囲を外れた場合、その時点でマッチは破棄されます。余計な計算をおこなわず高速に処理されます。マッチが成立した場合のみ、第2段階としてトレードモデリングに進みます。エントリーポイントは、パターン直後の新しいローソク足の始値として設定されます。

      //--- simulate a trade
      if(match)
        {
         local_stat[0] += 1;
         int open = i + PATTERN_SIZE;
         float4 bar = price[open];
         float entry = bar.x;
         float tp_val = tp[itp];
         float sl_val = sl[isl];
         float buy_tp  = entry + tp_val;
         float buy_sl  = entry - sl_val;
         float sell_tp = entry - tp_val;
         float sell_sl = entry + sl_val;
         bool buy_tp_hit  = 0, buy_sl_hit  = 0;
         bool sell_tp_hit = 0, sell_sl_hit = 0;
         for(int k = 0; k < horizon; k++)
           {
            bar = price[open + k];
            float high = bar.y;
            float low  = bar.z;
            // SL is checked first (worst-case)
            buy_sl_hit  |= (buy_tp_hit  == 0) & (buy_sl_hit  == 0) & (low  <= buy_sl);
            buy_tp_hit  |= (buy_tp_hit  == 0) & (buy_sl_hit  == 0) & (high >= buy_tp);
            sell_sl_hit |= (sell_tp_hit == 0) & (sell_sl_hit == 0) & (high >= sell_sl);
            sell_tp_hit |= (sell_tp_hit == 0) & (sell_sl_hit == 0) & (low  <= sell_tp);
            if((buy_tp_hit | buy_sl_hit) & (sell_tp_hit | sell_sl_hit))
               break;
           }
         // forced closing by time
         buy_sl_hit  |= (buy_tp_hit  == 0) & (buy_sl_hit  == 0) & (entry  > bar.w);
         buy_tp_hit  |= (buy_tp_hit  == 0) & (buy_sl_hit  == 0) & (entry  < bar.w);
         sell_sl_hit |= (sell_tp_hit == 0) & (sell_sl_hit == 0) & (entry  < bar.w);
         sell_tp_hit |= (sell_tp_hit == 0) & (sell_sl_hit == 0) & (entry  > bar.w);
         //---
         local_stat[1] += (int)buy_tp_hit;
         local_stat[2] += (int)buy_sl_hit;
         local_stat[3] += (int)sell_tp_hit;
         local_stat[4] += (int)sell_sl_hit;
        }
     }

次に、買いと売りの両方についてTPとSLのレベルが計算されます。ここで重要なのは、両サイドを同時にカウントしている点です。これにより計算量を削減し、市場挙動の全体像を得ることができます。

その後、時間的なホライゾン制約付きで履歴のフォワードパスが開始されます。各ステップでは次の点が確認されます。

• SLに到達したか
• TPに到達したか

SLのチェックが先におこなわれます。これは偶然ではなく、最悪ケースを仮定した意図的な設計です。このアプローチにより評価はより保守的になり、現実に近い結果になります。

TPまたはSLのいずれかに到達した時点で、そのポジションの結果は確定し、ループは終了します。せん。無駄な処理はおこなわれません。

TPもSLも到達しない場合は、時間制限によって強制決済されます。これも重要な要素です。ポジションを未決済のまま残さず、すべてのケースが必ず結果を持つように設計されています。

すべての結果はローカルのprivate local_stat配列に蓄積されます。ここが本質的なポイントです。この段階では同期処理はおこなわれません。各スレッドは独立して高速に処理をおこないます。

次に、全体設計の核心であるローカル集約に移ります。最初のBLOCK_SIZE個のスレッドがそれぞれの結果をbuf_statへ書き込みます。これはローカルメモリであり、高速かつワークグループ内で共有されます。

//--- write to 'local'
   if(lid < BLOCK_SIZE)
      for(int k = 0; k < STAT_DIM; k++)
         buf_stat[lid][k] = local_stat[k];
   barrier(CLK_LOCAL_MEM_FENCE);

その後、追加のパスが実行されます。これは、スレッド数がバッファサイズを超える場合にデータを圧縮する処理です。これは、任意のグループサイズを固定のリダクションウィンドウへ強制的に収束させる巧妙な方法です。

   for(int i = BLOCK_SIZE; i < total_loc; i += BLOCK_SIZE)
     {
      if(lid >= i && lid < (i + BLOCK_SIZE))
         for(int k = 0; k < STAT_DIM; k++)
            buf_stat[lid-i][k] += local_stat[k];
      barrier(CLK_LOCAL_MEM_FENCE);
     }

その後、クラシカルなリダクション処理がおこなわれます。ステップを段階的に縮小しながら、ペアごとの加算が実行されます。

//--- reduction
   for(int stride = BLOCK_SIZE / 2; stride > 0; stride >>= 1)
     {
      if(lid < stride)
        {
         for(int k = 0; k < STAT_DIM; k++)
           {
            buf_stat[lid][k] += buf_stat[lid + stride][k];
            buf_stat[lid + stride][k] = 0;
           }
        }
      barrier(CLK_LOCAL_MEM_FENCE);
     }

出力は1つのストリームで、グループ全体の集計統計情報が含まれています。最終段階として、結果を記録します。

//--- write the result
   if(lid == 0)
     {
      int idx = (itp * sl_count + isl) * STAT_DIM;
      int count = buf_stat[0][0];
      global_stats[idx + 0] = count;
      float norm = (count > 0) ? (1.0f / ((float)count)) : 0.0f;
      for(int k = 1; k < STAT_DIM; k++)
         global_stats[idx + k] = buf_stat[0][k] * norm;
     }
  }

ここで重要な変換が行われます。

• 総マッチ数はそのまま保持されます。
• 一方で、それ以外の値は確率へと変換されます。

計算が完了すると、出力は単なる生の数値ではなく、すぐに利用可能な統計情報になります。そこから次の情報が得られます。

• 類似した状況が過去に何回発生したか
• 買いにおいてテイクプロフィットが発動した頻度
• ストップロスが発動した頻度
• 売りがどのように推移したか
• どのTP/SLの組み合わせがより有利だったか

その後、メインプログラムが意思決定をおこないます。そこではパターンの見た目ではなく統計が参照されます。データが少なすぎる場合、そのシグナルは無視されます。十分なサンプルがある場合は、買いと売りの確率が比較されます。その上で、より適切なテイクプロフィット値とストップロスの値が選択されます。それらが揃って初めてポジションが開かれます。

これは重要な点です。この構造では、意思決定は推測や固定ルールに基づくものではありません。類似状況において市場が通常どのように振る舞ってきたか、その検証結果に基づいています。このアプローチは利益を保証するものではありませんが、システムとしてのデータ分析という考え方にははるかに整合的です。

テスト結果はこれを裏付けています。このシステムは完璧ではなく、非常に綺麗なエクイティカーブを示すわけでもありません。むしろそれは良い点とも言えます。ここにあるのは過剰適合したモデルではなく、通常の市場ノイズや連敗、ドローダウンにさらされる現実的なシンプル構造です。それでもなお、正の結果を示しています。

特に重要なのは、最終的な利益が中程度であることではなく、従来のインジケーターや事前定義されたパターンに依存せずに、統計的に意味のある取引機構を構築できた点です。この構造においてOpenCLは、こうした分析を時間的に現実的なものにするアクセラレータとして機能しています。

言い換えれば、ここでGPUが必要なのは戦略の収益性のためではありません。多数の過去の類似ケースを高速に検証し、市場データを測定可能な仮説へと変換するためです。

結論：OpenCLはローソク足パターン分析に有用です。なぜなら、多様な類似状況を高速に走査し、トレード結果を計算し、視覚的な推測ではなく統計に基づいて判断できるようにするためです。

結論

MQL5におけるOpenCLは、特定クラスのタスク向けツールとして捉えるべきです。その真価を発揮するのは、大量のデータ、繰り返し処理、そして計算を効率的に並列化できる状況です。それ以外の場合では、CPUの方が依然として合理的な選択肢となります。CPUはより単純で柔軟性が高く、小規模または並列化しにくい計算では、しばしばGPUより高速です。

本記事では、アーキテクチャ上の制約の理解から始まり、GPUへ計算をオフロードする実践的な構造に至るまでを一通り示しました。その中で明らかになったのは、結果を決定するのはGPUそのものの性能だけではなく、計算パイプラインの構成品質であるという点です。頻繁な初期化、冗長なデータ転送、小さすぎるタスクは、容易にGPUの利点を打ち消します。逆に、コンテキストを再利用し、大規模配列を扱い、データ受け渡しを慎重に設計した場合、GPUは本来の性能を発揮し始めます。

これは行列乗算の例で特に明確に示されました。十分なデータ量があれば、並列処理によってCPUに対して数倍の高速化が得られます。しかし、この構造の実用的価値は単なるベンチマークだけではありません。より実際の取引に近いタスクでは、GPUは大規模最適化、仮説検証、データ内の安定したパターン探索といった研究を高速化します。ここで初めて、GPUが単独で価値を持つのではなく、「計算規模を拡張するための手段」として有用であることが明確になります。

最も重要な結論は、あくまで実践的なものです。GPUへ処理を移したからといって、取引システムが自動的に収益化されるわけではなく、市場モデルそのものを置き換えることもできません。しかし、CPUではコストが高すぎる、あるいは時間がかかりすぎるような大量の組み合わせ探索や分析を可能にします。そこに本当の価値があります。奇跡を約束するのではなく、MQL5における、より広範かつ体系的な研究サイクルを実現すること。それがOpenCLの役割です。

記事で使用されているプログラム

#	名前	種類	詳細
1	PatternStats.mq5	EA	EAのテスト
2	PatternStats.cl	コードベース	OpenCLプログラムコードライブラリ