Multi-Tale of Lost Time
In the publication The Tale of Lost Time, user crea7or told how he refuted the Euler Hypothesis on a modern CPU.
I was also interested to know how the GPU will show itself, and I compared the single-threaded code with the multi-threaded for the CPU and the completely multi-threaded for the GPU, using the CUDA parallel computing architecture.
For a start, I parallelized the algorithm for the CPU. For this, the function is passed a range for iterating over the values of the outer loop a . Then the whole range of 1..N is broken into pieces and fed to the processor cores.
In my case, there were 4 cores, and I performed the calculations through a slightly modified thread pool, which I used as the basis here . Since there are no dependencies between the data, the speed increases almost in proportion to the number of cores.
')
For a GPU it is a little more difficult. The outer two cycles of searching ( a and b ) will be expanded, and in the function there will be only a search of two internal cycles ( c and d ). One can imagine that the program will be executed in parallel for all values a = 1..N and b = 1..N. In fact, of course, this is not the case; they will not be able to execute everything in parallel at the same time, but this is CUDA’s concern to parallelize the implementation as much as possible, this can be helped if the configuration of blocks and streams is properly configured.
The code for the GPU is quite straightforward and very similar to the CPU, only the binary search has to be done by hand.
I also implemented the optimizations suggested here .
The measurements for the CPU were done two times, since they turned out to be much more stable than the GPU, which I have already done seven and selected the best time. The variation for the GPU could be twofold, I explain this by the fact that the only video adapter in the system was used for CUDA calculations.
And this adrenaline is Profile-guided optimization .
For PGO, I cite only CPU time, because the GPU has changed little, and this is expected.
It can be seen that PGO was able to accelerate optimization # 1 by as much as 18%, for the rest the increase was more modest.
Forgot to mention a funny feature. At first I searched for the first solution, and set return immediately after printf. So it reduced productivity by an order of magnitude. When I began to look for all the solutions, the performance jumped sharply upwards.
For the CPU, try to expand the cycles, count two numbers at a time, use vector processor commands.
For the GPU to play with the order of calculations, try to save on registers, it is better to choose the parameters of the blocks and threads.
1. Writing on CUDA is easy.
2. Without tricks, the straight-line code on CUDA turned out to be 10 times faster than the CPU implementation.
3. In the number of crushed tasks, the GPU is much faster than the CPU for the same money.
4. The GPU takes a lot of time to "swing" and for very short tasks acceleration may not be.
5. A more complex algorithm may be faster for the CPU, but slower for the GPU.
6. Profile optimization speeds up code well and reduces file size simultaneously.
Traditionally you can touch the code on github
I was also interested to know how the GPU will show itself, and I compared the single-threaded code with the multi-threaded for the CPU and the completely multi-threaded for the GPU, using the CUDA parallel computing architecture.
Iron and compilers
CPU Core i5 4440
GPU GeForce GTX 770
MSVS 2015 update 3
tookit CUDA 8 .
Of course, the build was release and 64 bit, with optimization / 02 and / LTCG.
GPU GeForce GTX 770
MSVS 2015 update 3
tookit CUDA 8 .
Of course, the build was release and 64 bit, with optimization / 02 and / LTCG.
Disable video adapter reset
Since the calculations last more than two seconds, the video driver completes the CUDA program. To prevent this from happening, you need to specify a sufficient time before resetting through the registry key.
HKEY_LOCAL_MACHINE \ SYSTEM \ CurrentControlSet \ Control \ GraphicsDrivers \ TdrDelay and restart the computer.
HKEY_LOCAL_MACHINE \ SYSTEM \ CurrentControlSet \ Control \ GraphicsDrivers \ TdrDelay and restart the computer.
CPU
For a start, I parallelized the algorithm for the CPU. For this, the function is passed a range for iterating over the values of the outer loop a . Then the whole range of 1..N is broken into pieces and fed to the processor cores.
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"; } } } } } } In my case, there were 4 cores, and I performed the calculations through a slightly modified thread pool, which I used as the basis here . Since there are no dependencies between the data, the speed increases almost in proportion to the number of cores.
')
GPU
For a GPU it is a little more difficult. The outer two cycles of searching ( a and b ) will be expanded, and in the function there will be only a search of two internal cycles ( c and d ). One can imagine that the program will be executed in parallel for all values a = 1..N and b = 1..N. In fact, of course, this is not the case; they will not be able to execute everything in parallel at the same time, but this is CUDA’s concern to parallelize the implementation as much as possible, this can be helped if the configuration of blocks and streams is properly configured.
Blocks and threads
blocks_x is the enumeration range of the first cycle a .
But going through the second cycle b had to be made more difficult due to the video card restriction on the number of threads (1024), and put the counter simultaneously in threads and blocks_y :
a and b is calculated as:
This is definitely not the optimal configuration, as it turns out a lot of idle cycles. But I did not begin to adjust the values further.
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 is the enumeration range of the first cycle a .
But going through the second cycle b had to be made more difficult due to the video card restriction on the number of threads (1024), and put the counter simultaneously in threads and blocks_y :
a and b is calculated as:
const int a = blockIdx.x + 1; const int b = threadIdx.x + blockIdx.y * blockDim.x + 1; This is definitely not the optimal configuration, as it turns out a lot of idle cycles. But I did not begin to adjust the values further.
The code for the GPU is quite straightforward and very similar to the CPU, only the binary search has to be done by hand.
__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); } } } } I also implemented the optimizations suggested here .
Speed measurement
The measurements for the CPU were done two times, since they turned out to be much more stable than the GPU, which I have already done seven and selected the best time. The variation for the GPU could be twofold, I explain this by the fact that the only video adapter in the system was used for CUDA calculations.
| N | CPU time, ms | CPU (4 threads) time, ms | GPU time, ms | |
|---|---|---|---|---|
| @antonkrechetov | 100 | 58.6 | 16.7 | 14.8 |
| Base | 100 | 45.3 | 13.0 | 10.7 |
| Basic optimization # 1 | 100 | 6.3 | 2.1 | 12.7 |
| Basic Optimization # 2 | 100 | 1.4 | 0.7 | 0.8 |
| @antonkrechetov | 250 | 2999.7 | 782.9 | 119.0 |
| Base | 250 | 2055.6 | 550.1 | 90.9 |
| Basic optimization # 1 | 250 | 227.2 | 60.3 | 109.2 |
| Basic Optimization # 2 | 250 | 42.9 | 11.9 | 6.0 |
| @antonkrechetov | 500 | 72034.2 | 19344.1 | 1725.83 |
| Base | 500 | 38219.7 | 10200.8 | 976.7 |
| Basic optimization # 1 | 500 | 3725.1 | 926.5 | 1140.36 |
| Basic Optimization # 2 | 500 | 630.7 | 170.2 | 48.7 |
| @antonkrechetov | 750 | 396566.0 | 105003.0 | 11521.2 |
| Base | 750 | 218615.0 | 57639.2 | 5742.5 |
| Basic optimization # 1 | 750 | 19082.7 | 4736.8 | 6402.1 |
| Basic Optimization # 2 | 750 | 3272.0 | 846.9 | 222.9 |
| Basic Optimization # 2 | 1000 | 10204.4 | 2730.3 | 1041.9 |
| Basic Optimization # 2 | 1250 | 25133.1 | 6515.3 | 2445.5 |
| Basic Optimization # 2 | 1500 | 51940.1 | 14005.0 | 4895.2 |
And now give a little adrenaline to the CPU!
And this adrenaline is Profile-guided optimization .
For PGO, I cite only CPU time, because the GPU has changed little, and this is expected.
| N | CPU time, ms | CPU (4 threads) time, ms | CPU + PGO time, ms | CPU + PGO (4 threads) time, ms | |
|---|---|---|---|---|---|
| @antonkrechetov | 100 | 58.6 | 16.7 | 55.3 | 15.0 |
| Base | 100 | 45.3 | 13.0 | 42.2 | 12.1 |
| Basic optimization # 1 | 100 | 6.3 | 2.1 | 5.2 | 1.9 |
| Basic Optimization # 2 | 100 | 1.4 | 0.7 | 1.3 | 0.8 |
| @antonkrechetov | 250 | 2999.7 | 782.9 | 2966.1 | 774.1 |
| Base | 250 | 2055.6 | 550.1 | 2050.2 | 544.6 |
| Basic optimization # 1 | 250 | 227.2 | 60.3 | 200.0 | 53.2 |
| Basic Optimization # 2 | 250 | 42.9 | 11.9 | 40.4 | 11.4 |
| @antonkrechetov | 500 | 72034.2 | 19344.1 | 68662.8 | 17959.0 |
| Base | 500 | 38219.7 | 10200.8 | 38077.7 | 10034.0 |
| Basic optimization # 1 | 500 | 3725.1 | 926.5 | 3132.9 | 822.2 |
| Basic Optimization # 2 | 500 | 630.7 | 170.2 | 618.1 | 160.6 |
| @antonkrechetov | 750 | 396566.0 | 105003.0 | 404692.0 | 103602.0 |
| Base | 750 | 218615.0 | 57639.2 | 207975.0 | 54868.2 |
| Basic optimization # 1 | 750 | 19082.7 | 4736.8 | 15496.4 | 4082.3 |
| Basic Optimization # 2 | 750 | 3272.0 | 846.9 | 3093.8 | 812.7 |
| Basic Optimization # 2 | 1000 | 10204.4 | 2730.3 | 9781.6 | 2565.9 |
| Basic Optimization # 2 | 1250 | 25133.1 | 6515.3 | 23704.3 | 6244.1 |
| Basic Optimization # 2 | 1500 | 51940.1 | 14005.0 | 48717.5 | 12793.5 |
It can be seen that PGO was able to accelerate optimization # 1 by as much as 18%, for the rest the increase was more modest.
Dragons are found here
Forgot to mention a funny feature. At first I searched for the first solution, and set return immediately after printf. So it reduced productivity by an order of magnitude. When I began to look for all the solutions, the performance jumped sharply upwards.
__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; <- } } } } What else can you do
For the CPU, try to expand the cycles, count two numbers at a time, use vector processor commands.
For the GPU to play with the order of calculations, try to save on registers, it is better to choose the parameters of the blocks and threads.
findings
1. Writing on CUDA is easy.
2. Without tricks, the straight-line code on CUDA turned out to be 10 times faster than the CPU implementation.
3. In the number of crushed tasks, the GPU is much faster than the CPU for the same money.
4. The GPU takes a lot of time to "swing" and for very short tasks acceleration may not be.
5. A more complex algorithm may be faster for the CPU, but slower for the GPU.
6. Profile optimization speeds up code well and reduces file size simultaneously.
Sources
Traditionally you can touch the code on github
Source: https://habr.com/ru/post/318066/
All Articles