出版物
「The Tale of Lost Time」で、ユーザー
crea7orは、現代のCPUで
オイラー仮説をどのように
反fuしたかを説明しました。
GPUがどのように表示されるかを知りたいと思い、
CUDA並列コンピューティングアーキテクチャを使用して、シングルスレッドコードとCPUのマルチスレッドおよびGPUの完全マルチスレッドを比較しました。
アイロンとコンパイラCPU
コアi5 4440GPU
GeForce GTX 770MSVS 2015アップデート3Takeit 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計算に使用されたという事実によって説明できます。
そして、アドレナリンを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のコードに触れることができ
ます