RidgeRun CUDA Optimisation Guide/Empirical Experiments/Multi-threaded bounding test
|
|
Introduction
This page is the follow-up of Cuda Memory Benchmark, it adds multithreading to the testing of the different memory management modes. It's reduced to test traditional, managed, page-locked memory with and without copy call and CUDA mapped. [PENDING]
Testing Setup
Memory Management Methods
The program tested had the option to use each of the following memory management configurations:
- Traditional mode, using malloc to reserve the memory on host, then cudaMalloc to reserve it on the device, and then having to move the data between them with cudaMemcpy. Internally, the driver will allocate a non-pageable memory chunk, to copy the data there and after the copy, finally use the data on the device.
- Managed, using cudaMallocManaged and not having to manually copy the data and handle two different pointers.
- Non paging memory, using cudaMallocHost a chunk of page-locked memory can be reserved that can be used directly by the device since its non-pageable
- Non paging memory with discrete copy, using cudaMallocHost and a discrete call to cudaMemcpy, so its similar to the traditional model with different pointers one for host and another for device, but according to the NVIDIA docs on the mallocHost, the calls to cudaMemcpy are accelerated when using thid type of memory.
- CUDA mapped, using cudaHostAlloc to reserve memory that is page-locked and directly accessible to the device. There are different flags that can change the properties of the memory, in this case, the flags used were cudaHostAllocMapped and cudaHostAllocWriteCombined.
Platforms
- Discrete GPU: desktop pc with RTX 2070s, using CUDA 12 and Ubuntu 20.04.
- Jetson AGX Orin: using CUDA 11.4 and JP 5.0.2.
- Jetson Nano 4GB devkit: using CUDA 10.2 and JP 4.6.3
Program Structure
The program is divided into three main sections, one where the input memory is filled with data, the kernel worker threads, and the verify. The verify reads all the results and uses assert to verify them. Before every test, 10 iterations of the full process were done to warm up and avoid any initialization time penalty. After that, the average of 100 runs was obtained. Each of the sections can be seen in Figure 1.
Each kernel block can be seen in Figure 2. Each block has a semaphore to syncronize where there is data available on the input and has another that is raised at the end when the data is processed. The semaphores are shared between each block, in a chained manner, where the output semaphore is the input of the next block and so on. Also the kernel block has a condition that checks if the memory mode being used is "managed like" meaning that uses a single pointer to device and host memory and the driver handles the transfers on the back. If so it doesn't do the copy from host to device and back when done, otherwise it calls cudaMemcpy accordingly. And lastly the kernel itself it is called once for the IO bound case and 50 times for the processing bound case.
Used Data
The aim of the tests is to emulate a 4k RGBA frame so the results can be representative of the results on a real world media-handling software. To represent this data the following structure was used:
struct rgba_frame{ float r[SIZE_W*SIZE_H]; float g[SIZE_W*SIZE_H]; float b[SIZE_W*SIZE_H]; float a[SIZE_W*SIZE_H]; };
The macros are SIZE_W=3840 and SIZE_H=2160, the image size of a 4k frame.
Code
The kernel that was tested is:
Normalizer
int blockSize = 256; int numBlocks = ((SIZE_W*SIZE_H) + blockSize - 1) / blockSize; __global__ void normalize(int n, float *x, float *y) { int index = blockIdx.x * blockDim.x + threadIdx.x; int stride = blockDim.x * gridDim.x; for (int i = index; i < n; i += stride) y[i] = (x[i]*((MAX-MIN)/ABS_MAX))+MIN; } void *exec_kernel_cpy(void * arg){ struct sync_struct *args_k = (struct sync_struct *)arg; rgba_frame *d_frame_in = args_k->d_frame_in; rgba_frame *d_frame_out = args_k->d_frame_out; rgba_frame *frame_in = args_k->frame_in; rgba_frame *frame_out = args_k->frame_out; for (int i = 0; i < LOOP_CYCLES; i ++){ sem_wait(args_k->lock_in); if (i > WARM_UP_RUNS){ start_time(args_k); } if (d_frame_in != frame_in && mode != 4){ cudaMemcpy(d_frame_in, frame_in, sizeof(struct rgba_frame), cudaMemcpyHostToDevice); } for (int j = 0; j < KERNEL_CYCLES; j ++){ normalize<<<numBlocks, blockSize>>>(SIZE_W*SIZE_H, d_frame_in->r, d_frame_out->r); normalize<<<numBlocks, blockSize>>>(SIZE_W*SIZE_H, d_frame_in->g, d_frame_out->g); normalize<<<numBlocks, blockSize>>>(SIZE_W*SIZE_H, d_frame_in->b, d_frame_out->b); normalize<<<numBlocks, blockSize>>>(SIZE_W*SIZE_H, d_frame_in->a, d_frame_out->a); cudaDeviceSynchronize(); } if (frame_out != d_frame_out && mode != 4){ cudaMemcpy(frame_out, d_frame_out, sizeof(struct rgba_frame), cudaMemcpyDeviceToHost); } if (i > WARM_UP_RUNS){ stop_time(args_k); } sem_post(args_k->lock_out); } return NULL; }
This is the kernel that is bound to the worker threads, and containes two main cycles, one that is responsible to execute the number of loops to get the average and the inner one that uses the macro KERNEL_CYCLES. The value was changed between 1 and 50 to have a IO-bound case and processing-bound case, respectively. This can be seen on the figures with the label normalizer 1x and normalizer 50x, respectively.
Apart from this, the code has two sections an initial section and a end section. The initial section takes the array and fills it with 1s. It also contains the cycle responsible for the average.
void *fill_array(void * arg){ struct sync_struct *args_fill = (struct sync_struct *)arg; for (int i = 0; i < LOOP_CYCLES; i ++){ sem_wait(args_fill->lock_in); if (i > WARM_UP_RUNS){ start_time(args_fill); } for (int j = 0; j < SIZE_W*SIZE_H; j++) { args_fill->frame_out->r[j] = 1.0f; } for (int j = 0; j < SIZE_W*SIZE_H; j++) { args_fill->frame_out->g[j] = 1.0f; } for (int j = 0; j < SIZE_W*SIZE_H; j++) { args_fill->frame_out->b[j] = 1.0f; } for (int j = 0; j < SIZE_W*SIZE_H; j++) { args_fill->frame_out->a[j] = 1.0f; } if (i > WARM_UP_RUNS){ stop_time(args_fill); } sem_post(args_fill->lock_out); } return NULL; }
The end section is where the output is read and at the same time, the results are checked to verify that the process is behaving as expected.
void *verify_results(void * arg){ struct sync_struct *args_verf = (struct sync_struct *)arg; float ref = 1.0f; for (int i = 0; i < STAGES-1; i++){ ref = (ref*((MAX-MIN)/ABS_MAX))+MIN; } for (int i = 0; i < LOOP_CYCLES; i ++){ sem_wait(args_verf->lock_in); if (i > WARM_UP_RUNS){ start_time(args_verf); } for (int j = 0; j < SIZE_W*SIZE_H; j++) { assert(args_verf->frame_in->r[j] == ref); } for (int j = 0; j < SIZE_W*SIZE_H; j++) { assert(args_verf->frame_in->g[j] == ref); } for (int j = 0; j < SIZE_W*SIZE_H; j++) { assert(args_verf->frame_in->b[j] == ref); } for (int j = 0; j < SIZE_W*SIZE_H; j++) { assert(args_verf->frame_in->a[j] == ref); } if (i > WARM_UP_RUNS){ stop_time(args_verf); } sem_post(args_verf->lock_out); } return NULL; }
Each section was measured using different timers, and the results were added to get the total times. As can be seen from the code pieces each worker has a sync_struct associated, this piece most notably, holds the semaphores for each and the times for each, amongs other necessary values.
Results
Discrete GPU
Kernel Execution Time
In Figure 3 the fastest mode is CUDA mapped, at 11.85ms on average, followed by pinned memory without copy and managed on the IO-bound case, but on the processing bound case, CUDA mapped and pinned memory without copy, suffer an increase of around 48 times for both. And in this scenario, the fastes mode is managed, and Table 1 sheds some light into the reason.
Case | Worker | Time avg(ms) |
---|---|---|
IO-Bound | W1 | 30.466015 |
W2 | 0.761422 | |
W3 | 0.757867 | |
W4 | 0.757999 | |
W5 | 29.853945 | |
Processing-Bound | W1 | 64.937218 |
W2 | 36.428116 | |
W3 | 36.336861 | |
W4 | 29.853945 | |
W5 | 64.063713 |
It can be seen that in both scenarios the CUDA runtime identifies the chain of buffers and speeds up the data transfers which results on a considerable 40 times, time reduction on the IO-bound case and 2 times, on the Processing-Bound on the inner worker threads. However it can also be seen that there is a time penalty on the end thread and initial thread as it has to fetch the data from host memory. No other memory mode showed this behavior.
Full Execution Times
On table 2, it can be seen that the best overall is hostMalloc, however it's between variance of the traditional mode. Also the worst overall with a verify time of 58 times more that the best is CUDA mapped.
Memory mode | Fill time avg(ms) | Verify time avg(ms) |
---|---|---|
Traditional | 63.391980 | 51.470143 |
Managed | 62.733280 | 71.642704 |
HostMalloc | 61.780409 | 51.822929 |
HostMalloc W Cpy | 64.2764265 | 51.49177 |
Pinned | 72.431336 | 3,001.322266 |
When we take the three times and combine them, to get the total execution time, as shown in Figure 4. We see that in the case of the discrete GPU, the best performing for the IO-Bound case is HostMalloc without discrete copy, and for the Processing-Bound case, the best is Managed memory, since it has that edge on the worker to worker transfers.
In general, it seems that in IO-bound cases, it can yield benefits using memory reserved with hostMalloc and not doing the manual copy, but on a processing-bound scenario, the dicrete call to copy is needed. Overall we have slower performance with managed memory and the slowest is with pinned or zero-copy memory.
Jetson Nano
Kernel Execution Time
On Figure 5, we can see that the mode that performs better is CUDA mapped, and the next is pinned memory without copy, followed by managed, which is the same trend that the dGPU results had on the IO-bound case. But on the Processing-Bound case the behavior is different, where pinned memory with copy performs the best and it's followed by the same without the copy. Also worth pointing out is that the time difference between kernel workers, is not present on the results. On the Jetson Nano the results for each kernel are close to each other.
Full Execution Times
On Table
Memory mode | Fill time avg(ms) | Verify time avg(ms) |
---|---|---|
Traditional | 352.960114 | 190.085793 |
Managed | 398.220825 | 241.442940 |
HostMalloc | 351.453018 | 1,243.674988 |
HostMalloc W Cpy | 351.932023 | 1,244.144898 |
Pinned | 351.853775 | 1,243.925659 |
Figure 5, shows that the Jetson Nano has a different trend, where the managed actually performs well. We have the same behavior as the discrete GPU for the pinned or zero copy. But the hostMalloc performs slowly, compared to the discrete GPU results.
Overall it seems that the managed performs better on Jetson Nano than on the discrete GPU. In this case, it does not make sense to use pinned or zero copy nor memory reserved with hostMalloc. In the IO-bound case, managed memory can perform better than traditional, but on a processing-bound program, the traditional performs best.
Jetson AGX Orin(pending)
Kernel execution time
In kernel execution times, Figure 6, there is a clear time reduction when using hostMalloc, where it performs better than traditional memory management. It can be seen that with managed memory there is a bit more gain to be had at around 2ms less than either of them.
For the Jetson AGX Orin, we have results that look more like the discrete GPU results, but the main difference is that using memory with hostMalloc does achieve better results always.
Full Execution Times
As for the Jetson AGX Orin, Table 3, the results from the fill and verify operations, show that there is a similar trend as the discrete GPU, where the managed performs slower compared to the rest, but as for the pinned, there is not as much time increase compared to the discrete GPU, since its around 3 times more, but still, it performs the worst.
Memory mode | Fill time avg(ms) | Verify time avg(ms) |
---|---|---|
Traditional | 96.3693465 | 93.119007 |
Managed | 141.320404 | 90.6683695 |
Managed & prefetch | 140.552361 | 90.991768 |
Managed & advice as GPU | 141.023903 | 91.376148 |
Managed & advice as CPU | 141.235008 | 91.276241 |
Managed & prefetch & advice as GPU | 141.092499 | 91.323822 |
Managed & prefetch & advice as CPU | 140.9354705 | 91.141693 |
HostMalloc | 96.887695 | 99.5928345 |
HostMalloc W Cpy | 96.2297175 | 98.4038125 |
Pinned | 96.558895 | 792.3588565 |
When looking at the full execution times, Figure 7, there is a different behavior than the Jetson Nano, but it's similar to the discrete GPU. Where the managed performs notably slower overall, and the hostMalloc performs better.
In the case of the Jetson Orin AGX, there is one memory management mode that performs better without regarding if its an IO o processing bound scenario, that being the memory reserved with hostMalloc and without the need of handling the discrete transfers, compared to the discrete GPU.
Resource Usage Jetson(pending)
In both Jetson targets, tegrastats was used to monitor the resource utilization, mainly the CPU and GPU usage and the used memory. Upon inspection, there is virtually no difference from run to run. Where the different memory management tests, used the same amount of memory. As for the general system usage, there is also nothing worthy of attention.
Conclusions(pending)
We don't have a definitive management mode that performs best in all cases and all devices, but we can see that in different use cases and devices, one can perform better than the other. However, if you are looking for consistency and control, the traditional memory model is the way to go. But if you need to have the best execution times, we have some points that might help:
- On a discrete GPU, use the hostMalloc memory model, but remember to use manual transfers when in a processing bound case.
- On Jetson Nano, on an IO-bound scenario, use managed memory, otherwise, use the traditional memory model.
- On Jetson AGX Orin we do have a one for all, in this case, use hostMalloc, this performs the best regardless, and with the bonus of not having to handle dual pointers for device and host memory.
RidgeRun Resources | |||||
Contact Us
|