👨‍🌾 🏵️ 👂🏽 失われた時間のマルチスレッド物語 🔋 ☔️ ✋

出版物「The Tale of Lost Time」で、ユーザーcrea7orは、現代のCPUでオイラー仮説をどのように反fuしたかを説明しました。

GPUがどのように表示されるかを知りたいと思い、 CUDA並列コンピューティングアーキテクチャを使用して、シングルスレッドコードとCPUのマルチスレッドおよびGPUの完全マルチスレッドを比較しました。

アイロンとコンパイラ

CPU コアi5 4440

GPU GeForce GTX 770

MSVS 2015アップデート3

Takeit CUDA 8 。

もちろん、ビルドはリリースおよび64ビットで、最適化は/ 02および/ LTCGでした。

ビデオアダプターのリセットを無効にする

計算には2秒以上かかるため、ビデオドライバーはCUDAプログラムを完了します。これを防ぐには、レジストリキーでリセットする前に十分な時間を指定する必要があります

HKEY_LOCAL_MACHINE \ SYSTEM \ CurrentControlSet \ Control \ GraphicsDrivers \ TdrDelayとコンピューターを再起動します。

CPU

まず、CPUのアルゴリズムを並列化しました。これを行うには、範囲を関数に渡し、外側のループaの値を反復処理します。次に、範囲1..N全体が断片に分割され、プロセッサコアに供給されます。

void Algorithm_1(const uint32_t N, const std::vector<uint64_t>& powers, const uint32_t from, const uint32_t to) { for (uint32_t a = from; a < to; ++a) { for (uint32_t b = 1; b < a; ++b) { for (uint32_t c = 1; c < b; ++c) { for (uint32_t d = 1; d < c; ++d) { const uint64_t sum = powers[a] + powers[b] + powers[c] + powers[d]; if (std::binary_search(std::begin(powers), std::end(powers), sum)) { const auto it = std::lower_bound(std::begin(powers), std::end(powers), sum); uint32_t e = static_cast<uint32_t>(std::distance(std::begin(powers), it)); std::cout << a << " " << b << " " << c << " " << d << " " << e << "\n"; } } } } } }

私の場合、4つのコアがあり、ここで基礎として取ったわずかに変更されたスレッドプールを介して計算を実行しました。データ間に依存関係がないため、速度はコアの数にほぼ比例して増加します。

GPU

GPUの場合、少し複雑になります。外側の2つの列挙ループ（ aおよびb ）がデプロイされ、関数は内側の2つのループ（ cおよびd ）を列挙するだけで済みます。プログラムは、 a = 1..Nおよびb = 1..Nのすべての値に対して並列に実行されると想像できます。実際、これはそうではありません。それらはまだ並列に実行できませんが、実行を可能な限り並列化することがCUDAの関心事です。これは、ブロックとフローの構成が正しく構成されている場合に役立ちます。

ブロックとストリーム

  int blocks_x = N; int threads = std::min(1024, static_cast<int>(N)); int blocks_y = (N + threads - 1) / threads; dim3 grid(blocks_x, blocks_y); NaiveGPUKernel<<<grid, threads>>>(N);

blocks_xは、最初のループaの列挙範囲です。

しかし、2番目のサイクルbの検索は、ビデオカードのスレッド数の制限（1024）により困難になり、 threadsとblocks_yに同時にカウンターを配置する必要がありました 。

aとbは次のように計算されます。

  const int a = blockIdx.x + 1; const int b = threadIdx.x + blockIdx.y * blockDim.x + 1;

多くのアイドルサイクルが発生するため、これは間違いなく最適な構成ではありません。しかし、値をそれ以上調整しませんでした。

GPUのコードは非常に単純で、CPUに非常に似ています。バイナリ検索のみを手で行う必要があります。

 __constant__ uint64_t gpu_powers[8192]; inline __device__ int gpu_binary_search(const uint32_t N, const uint64_t search) { uint32_t l = 0; uint32_t r = elements_count - 1; uint32_t m; while (l <= r) { m = (l + r) / 2; if (search < gpu_powers[m]) r = m - 1; else if (search > gpu_powers[m]) l = m + 1; else return l; } return -1; } __global__ void NaiveGPUKernel(const uint32_t N) { const int a = blockIdx.x + 1; const int b = threadIdx.x + blockIdx.y * blockDim.x + 1; if (b >= a) return; for (uint32_t c = 1; c < b; ++c) { for (uint32_t d = 1; d < c; ++d) { const uint64_t sum = gpu_powers[a] + gpu_powers[b] + gpu_powers[c] + gpu_powers[d]; const auto s = gpu_binary_search(N, sum); if (s > 0) { printf("%d %d %d %d %d\n", a, b, c, d, s); } } } }

また、ここで提案されている最適化も実装しました。

速度測定

CPUの測定は、GPUよりもはるかに安定していることが判明したため、2回行われました。これは、一度に7回行い、最適な時間を選択しました。 GPUの普及は2倍になる可能性があります。これは、システム内の唯一のビデオアダプターがCUDA計算に使用されたという事実によって説明できます。

	N	CPU時間、ミリ秒	CPU（4スレッド）時間、ms	GPU時間、ミリ秒
しょーた	100	58.6	16.7	14.8
ベーシック	100	45.3	13.0	10.7
基本最適化＃1	100	6.3	2.1	12.7
ベース最適化＃2	100	1.4	0.7	0.8
しょーた	250	2999.7	782.9	119.0
ベーシック	250	2055.6	550.1	90.9
基本最適化＃1	250	227.2	60.3	109.2
ベース最適化＃2	250	42.9	11.9	6.0
しょーた	500	72034.2	19344.1	1725.83
ベーシック	500	38219.7	10200.8	976.7
基本最適化＃1	500	3725.1	926.5	1140.36
ベース最適化＃2	500	630.7	170.2	48.7
しょーた	750	396566.0	105003.0	11521.2
ベーシック	750	218615.0	57639.2	5742.5
基本最適化＃1	750	19082.7	4736.8	6402.1
ベース最適化＃2	750	3272.0	846.9	222.9
ベース最適化＃2	1000	10204.4	2730.3	1041.9
ベース最適化＃2	1250	25133.1	6515.3	2445.5
ベース最適化＃2	1500	51940.1	14005.0	4895.2

そして、アドレナリンをCPUに注入しましょう！

そして、そのアドレナリンはプロファイルに基づく最適化になります。

PGOの場合、GPUがほとんど変化していないため、CPU時間のみを指定します。これは予想されることです。

	N	CPU 時間、ミリ秒	CPU（4スレッド）時間、ミリ秒	CPU + PGO 時間、ミリ秒	CPU + PGO （4スレッド）時間、ミリ秒
しょーた	100	58.6	16.7	55.3	15.0
ベーシック	100	45.3	13.0	42.2	12.1
基本最適化＃1	100	6.3	2.1	5.2	1.9
ベース最適化＃2	100	1.4	0.7	1.3	0.8
しょーた	250	2999.7	782.9	2966.1	774.1
ベーシック	250	2055.6	550.1	2050.2	544.6
基本最適化＃1	250	227.2	60.3	200.0	53.2
ベース最適化＃2	250	42.9	11.9	40.4	11.4
しょーた	500	72034.2	19344.1	68662.8	17959.0
ベーシック	500	38219.7	10200.8	38077.7	10034.0
基本最適化＃1	500	3725.1	926.5	3132.9	822.2
ベース最適化＃2	500	630.7	170.2	618.1	160.6
しょーた	750	396566.0	105003.0	404692.0	103602.0
ベーシック	750	218615.0	57639.2	207975.0	54868.2
基本最適化＃1	750	19082.7	4736.8	15496.4	4082.3
ベース最適化＃2	750	3272.0	846.9	3093.8	812.7
ベース最適化＃2	1000	10204.4	2730.3	9781.6	2565.9
ベース最適化＃2	1250	25133.1	6515.3	23704.3	6244.1
ベース最適化＃2	1500	51940.1	14005.0	48717.5	12793.5

PGOは最適化＃1を最大18％加速することができたことがわかりますが、それ以外の場合、増加はより緩やかでした。

ドラゴンがいます

楽しい機能について言及するのを忘れました。最初に最初の解決策を検索し、printfの直後にreturnを設定しました。そのため、パフォーマンスが一桁低下しました。すべてのソリューションを探し始めたとき、生産性は急激に上がりました。

 __global__ void NaiveGPUKernel(const uint32_t N) { const int a = blockIdx.x + 1; const int b = threadIdx.x + blockIdx.y * blockDim.x + 1; if (b >= a) return; for (uint32_t c = 1; c < b; ++c) { for (uint32_t d = 1; d < c; ++d) { const uint64_t sum = gpu_powers[a] + gpu_powers[b] + gpu_powers[c] + gpu_powers[d]; const auto s = gpu_binary_search(N, sum); if (s > 0) { printf("%d %d %d %d %d\n", a, b, c, d, s); return; <-      } } } }

さらにできること

CPUについては、サイクルを拡張し、一度に2つの数値を読み取り、ベクトルプロセッサ命令を使用してください。

GPUの場合、計算の順序で再生し、レジスタに保存してみてください。ブロックとフローのパラメーターを選択することをお勧めします。

結論

1. CUDAでの書き込みは簡単です。

2.トリックなしで、簡単なCUDAコードは実装CPUの10倍高速であることが判明しました。

3.数値計算タスクでは、同じ費用でGPUはCPUよりもはるかに高速です。

4. GPUは「スイング」するのに多くの時間を必要とし、非常に短いタスクの場合、加速はそうでない場合があります。

5.より複雑なアルゴリズムは、CPUでは高速ですが、GPUでは低速です。

6.プロファイルの最適化は、コードを高速化し、同時にファイルサイズを削減します。

ソースコード

従来、 GitHubのコードに触れることができます

失われた時間のマルチスレッド物語

CPU

GPU

速度測定

そして、アドレナリンをCPUに注入しましょう！

ドラゴンがいます

さらにできること

結論

ソースコード

More articles: