EXSWHTEC-249 - Implement Performance Tests for Memcpy APIs

[milos-mozetic]

Implement microbenchmarks for the Memcpy APIs.

The following should be covered:

• Synchronous and asynchronous variations of the APIs
• Small, medium, and large memory sizes
• Different copy directions (Host to Host, Host to Device, Device to Device, Device to Host)
• 1D, 2D, and 3D variations of the APIs
• From/to Symbol variations of the APIs
• From/to Array variations of the APIs

Depends on: #117

[chrispaquot]

Overall comments here:
I don't see loops of memcpy for example. Usually we like to loop and average out
I also don't see warm up calls. If you benchmark hipMemcpy, you should call it once at least before timing it for good measure.

[milos-mozetic]

Overall comments here: I don't see loops of memcpy for example. Usually we like to loop and average out I also don't see warm up calls. If you benchmark hipMemcpy, you should call it once at least before timing it for good measure.

Hello @chrispaquot, please check this file here: https://github.com/ROCm-Developer-Tools/hip-tests/pull/119/files#diff-5e0a673157e5c8f6a3c1cd2490e0039de1b789e42040f6f739b85575ca4c635a. You can see that within the Run method of a Benchmark class, the API that is being benchmarked is called several times for warmup and afterwards several times in the loop, which is used for timing and time metrics calculation.
This number of warmups and loop iterations used for benchmark metrics can be configured using command line arguments, as requested from AMD. Please check this file here: https://github.com/ROCm-Developer-Tools/hip-tests/pull/119/files#diff-c76ce2197e6cd6361be77f86ff31e6f31e7fd53df6bb88112d6f551aff7f8973.

[chrispaquot]

OK that sounds good with one more question: Iterations is going to be the same for all test cases then? Warmup would most likely be 1 but iterations can change from a test case to another.

[milos-mozetic]

OK that sounds good with one more question: Iterations is going to be the same for all test cases then? Warmup would most likely be 1 but iterations can change from a test case to another.

Yes, the number of iterations will be the same for all test cases within one logical group/module that will be executed. The request was to configure the number of warmups/iterations through command line arguments. You can configure specific number of warmups/iterations for each test case using Configure method from Benchmark class.

[gandryey]

There is a bit of problem. The current warm-up/iterations will run all allocations and destructions also. Warmup means it touches all resources and executes exactly the same commands, used in the final measurement. "Iterations" means the execution of the same commands in a loop multiple times with the same resources, used in a warm-up.

      TIMED_SECTION_STREAM(kTimerTypeEvent, stream) {
        HIP_CHECK(hipMemcpy2DAsync(host_allocation.ptr(), device_allocation.width(),
                                   device_allocation.ptr(), device_allocation.pitch(),
                                   device_allocation.width(), device_allocation.height(),
                                   hipMemcpyDeviceToHost, stream));
      }

[milos-mozetic]

There is a bit of problem. The current warm-up/iterations will run all allocations and destructions also. Warmup means it touches all resources and executes exactly the same commands, used in the final measurement. "Iterations" means the execution of the same commands in a loop multiple times with the same resources, used in a warm-up.

      TIMED_SECTION_STREAM(kTimerTypeEvent, stream) {
        HIP_CHECK(hipMemcpy2DAsync(host_allocation.ptr(), device_allocation.width(),
                                   device_allocation.ptr(), device_allocation.pitch(),
                                   device_allocation.width(), device_allocation.height(),
                                   hipMemcpyDeviceToHost, stream));
      }

Review comment addressed within 9ccdf02 and 96bb38d. Allocation guards are moved outside of warmups/iterations, so that the same commands are executed on the same resources. @gandryey, can you please review changes?

[rakesroy]

PR has been merged into develop branch via commit e3bac85.