Using hipStream to improve performance

whyb · whyb · commit 72d65c15ae35 · 2025-04-07T23:17:03.000+08:00
diff --git a/FastChwHwcConverterROCm.hpp b/FastChwHwcConverterROCm.hpp
@@ -552,26 +552,26 @@ inline void hwc2chw_rocm(
     const size_t input_size = pixel_size * sizeof(uint8_t);
     const size_t output_size = pixel_size * sizeof(float);
 
-    void* rocm_input_memory = nullptr;
-    void* rocm_output_memory = nullptr;
+    hipDeviceptr_t rocm_input_memory = 0;
+    hipDeviceptr_t rocm_output_memory = 0;
 
     // Allocate host-pinned memory
-    hipError_t hipRes0 = hipHostMalloc(&rocm_input_memory, input_size, 0);
-    hipError_t hipRes1 = hipHostMalloc(&rocm_output_memory, output_size, 0);
+    hipError_t hipRes0 = hipMallocAsync(&rocm_input_memory, input_size, rocmstream);
+    hipError_t hipRes1 = hipMallocAsync(&rocm_output_memory, output_size, rocmstream);
 
     if (hipRes0 != 0 || hipRes1 != 0) {
-        hipHostFree(rocm_input_memory);
-        hipHostFree(rocm_output_memory);
+        hipFreeAsync(rocm_input_memory, rocmstream);
+        hipFreeAsync(rocm_output_memory, rocmstream);
         hwc2chw<uint8_t, float>(h, w, c, src, dst, alpha);
         return;
     }
 
     // Copy host memory to device memory
-    hipError_t hipRes2 = hipMemcpyAsync(rocm_input_memory, src, input_size, hipMemcpyKind::hipMemcpyHostToDevice, rocmstream);
+    hipError_t hipRes2 = hipMemcpyHtoDAsync(rocm_input_memory, src, input_size, rocmstream);
 
     if (hipRes2 != 0) {
-        hipHostFree(rocm_input_memory);
-        hipHostFree(rocm_output_memory);
+        hipFreeAsync(rocm_input_memory, rocmstream);
+        hipFreeAsync(rocm_output_memory, rocmstream);
         hwc2chw<uint8_t, float>(h, w, c, src, dst, alpha);
         return;
     }
@@ -593,25 +593,25 @@ inline void hwc2chw_rocm(
         0, rocmstream, args, nullptr);
 
     if (hipRes3 != 0) {
-        hipHostFree(rocm_input_memory);
-        hipHostFree(rocm_output_memory);
+        hipFreeAsync(rocm_input_memory, rocmstream);
+        hipFreeAsync(rocm_output_memory, rocmstream);
         hwc2chw<uint8_t, float>(h, w, c, src, dst, alpha);
         return;
     }
 
     // Copy device memory to host memory
-    hipError_t hipRes5 = hipMemcpyAsync(dst, rocm_output_memory, output_size, hipMemcpyKind::hipMemcpyDeviceToHost, rocmstream);
+    hipError_t hipRes5 = hipMemcpyDtoHAsync(dst, rocm_output_memory, output_size, rocmstream);
 
     if (hipRes5 != 0) {
-        hipHostFree(rocm_input_memory);
-        hipHostFree(rocm_output_memory);
+        hipFreeAsync(rocm_input_memory, rocmstream);
+        hipFreeAsync(rocm_output_memory, rocmstream);
         hwc2chw<uint8_t, float>(h, w, c, src, dst, alpha);
         return;
     }
 
     // Free memory
-    hipHostFree(rocm_input_memory);
-    hipHostFree(rocm_output_memory);
+    hipFreeAsync(rocm_input_memory, rocmstream);
+    hipFreeAsync(rocm_output_memory, rocmstream);
 
     // Stream synchronization
     hipError_t hipRes4 = hipStreamSynchronize(rocmstream);
@@ -645,25 +645,25 @@ inline void chw2hwc_rocm(
     const size_t pixel_size = h * w * c;
     size_t input_size = pixel_size * sizeof(float);
     size_t output_size = pixel_size * sizeof(uint8_t);
-    void* rocm_input_memory = nullptr;
-    void* rocm_output_memory = nullptr;
+    hipDeviceptr_t rocm_input_memory = 0;
+    hipDeviceptr_t rocm_output_memory = 0;
 
     // Allocate device memory
-    hipError_t hipRes0 = hipHostMalloc(&rocm_input_memory, input_size, 0);
-    hipError_t hipRes1 = hipHostMalloc(&rocm_output_memory, output_size, 0);
+    hipError_t hipRes0 = hipMallocAsync(&rocm_input_memory, input_size, rocmstream);
+    hipError_t hipRes1 = hipMallocAsync(&rocm_output_memory, output_size, rocmstream);
 
     if (hipRes0 != 0 || hipRes1 != 0) {
-        hipHostFree(rocm_input_memory);
-        hipHostFree(rocm_output_memory);
+        hipFreeAsync(rocm_input_memory, rocmstream);
+        hipFreeAsync(rocm_output_memory, rocmstream);
         chw2hwc<float, uint8_t>(h, w, c, src, dst, alpha); return;
     }
 
     // Copy host memory to device memory
-    hipError_t hipRes2 = hipMemcpyAsync(rocm_input_memory, src, input_size, hipMemcpyKind::hipMemcpyHostToDevice, rocmstream);
+    hipError_t hipRes2 = hipMemcpyHtoDAsync(rocm_input_memory, src, input_size, rocmstream);
 
     if (hipRes2 != 0) {
-        hipHostFree(rocm_input_memory);
-        hipHostFree(rocm_output_memory);
+        hipFreeAsync(rocm_input_memory, rocmstream);
+        hipFreeAsync(rocm_output_memory, rocmstream);
         chw2hwc<float, uint8_t>(h, w, c, src, dst, alpha); return;
     }
 
@@ -686,23 +686,23 @@ inline void chw2hwc_rocm(
         0, rocmstream, args, nullptr);
 
     if (hipRes3 != 0) {
-        hipHostFree(rocm_input_memory);
-        hipHostFree(rocm_output_memory);
+        hipFreeAsync(rocm_input_memory, rocmstream);
+        hipFreeAsync(rocm_output_memory, rocmstream);
         chw2hwc<float, uint8_t>(h, w, c, src, dst, alpha); return;
     }
 
     // Copy device memory to host memory
-    hipError_t hipRes5 = hipMemcpyAsync(dst, rocm_output_memory, output_size, hipMemcpyKind::hipMemcpyDeviceToHost, rocmstream);
+    hipError_t hipRes5 = hipMemcpyDtoHAsync(dst, rocm_output_memory, output_size, rocmstream);
 
     if (hipRes5 != 0) {
-        hipHostFree(rocm_input_memory);
-        hipHostFree(rocm_output_memory);
+        hipFreeAsync(rocm_input_memory, rocmstream);
+        hipFreeAsync(rocm_output_memory, rocmstream);
         chw2hwc<float, uint8_t>(h, w, c, src, dst, alpha); return;
     }
 
     // Free memory
-    hipHostFree(rocm_input_memory);
-    hipHostFree(rocm_output_memory);
+    hipFreeAsync(rocm_input_memory, rocmstream);
+    hipFreeAsync(rocm_output_memory, rocmstream);
 
     hipError_t hipRes4 = hipStreamSynchronize(rocmstream);
 
@@ -725,7 +725,7 @@ inline void chw2hwc_rocm(
  */
 inline void hwc2chw_rocm(
     const size_t h, const size_t w, const size_t c,
-    void* src, void* dst,
+    hipDeviceptr_t src, hipDeviceptr_t dst,
     const float alpha = 1.f / 255.f) {
 
     const size_t pixel_size = h * w * c;
@@ -769,7 +769,7 @@ inline void hwc2chw_rocm(
  */
 inline void chw2hwc_rocm(
     const size_t c, const size_t h, const size_t w,
-    void* src, void* dst,
+    hipDeviceptr_t src, hipDeviceptr_t dst,
     const uint8_t alpha = 255.0f) {
     
     const unsigned int blockDimX = 32, blockDimY = 32, blockDimZ = 1;
diff --git a/README.md b/README.md
@@ -338,33 +338,33 @@ The table below shows the benchmark performance timing for different image dimen
 
     GPU(ROCm): AMD Radeon RX 6900 XT
 
-|       |        |         |    CPU   |    CPU   |CPU(OpenMP)|CPU(OpenMP)|   CUDA   |   CUDA   |   ROCm   |   ROCm  |
-|-------|--------|---------|----------|----------|-----------|-----------|----------|----------|----------|---------|
-|   W  |   H  |  C | hwc2chw  | chw2hwc  | hwc2chw   | chw2hwc   | hwc2chw  | chw2hwc  | hwc2chw  | chw2hwc |
-| 426  | 240  | 1  | 0.097ms  | 0.110ms  | 0.113ms   | 0.030ms   | 0.006ms  | 0.006ms  | 0.119ms  | 0.052ms |
-| 426  | 240  | 3  | 0.331ms  | 0.314ms  | 0.061ms   | 0.068ms   | 0.008ms  | 0.007ms  | 0.118ms  | 0.055ms |
-| 426  | 240  | 4  | 0.439ms  | 0.415ms  | 0.082ms   | 0.082ms   | 0.010ms  | 0.010ms  | 0.113ms  | 0.057ms |
-| 640  | 360  | 1  | 0.217ms  | 0.236ms  | 0.048ms   | 0.052ms   | 0.010ms  | 0.010ms  | 0.111ms  | 0.051ms |
-| 640  | 360  | 3  | 0.743ms  | 0.705ms  | 0.147ms   | 0.140ms   | 0.013ms  | 0.015ms  | 0.129ms  | 0.055ms |
-| 640  | 360  | 4  | 0.881ms  | 0.921ms  | 0.219ms   | 0.203ms   | 0.015ms  | 0.016ms  | 0.121ms  | 0.057ms |
-| 854  | 480  | 1  | 0.393ms  | 0.415ms  | 0.094ms   | 0.089ms   | 0.009ms  | 0.008ms  | 0.282ms  | 0.050ms |
-| 854  | 480  | 3  | 1.328ms  | 1.269ms  | 0.250ms   | 0.232ms   | 0.019ms  | 0.020ms  | 0.325ms  | 0.059ms |
-| 854  | 480  | 4  | 1.717ms  | 1.670ms  | 0.263ms   | 0.262ms   | 0.024ms  | 0.025ms  | 0.334ms  | 0.065ms |
-| 1280 | 720  | 1  | 0.873ms  | 0.937ms  | 0.130ms   | 0.180ms   | 0.019ms  | 0.016ms  | 0.301ms  | 0.053ms |
-| 1280 | 720  | 3  | 2.877ms  | 2.828ms  | 0.449ms   | 0.457ms   | 0.039ms  | 0.039ms  | 0.390ms  | 0.062ms |
-| 1280 | 720  | 4  | 3.558ms  | 3.848ms  | 0.719ms   | 0.616ms   | 0.049ms  | 0.051ms  | 0.336ms  | 0.063ms |
-| 1920 | 1080 | 1  | 1.949ms  | 2.136ms  | 0.374ms   | 0.342ms   | 0.036ms  | 0.032ms  | 0.616ms  | 0.055ms |
-| 1920 | 1080 | 3  | 6.587ms  | 6.469ms  | 1.000ms   | 0.672ms   | 0.081ms  | 0.084ms  | 0.827ms  | 0.065ms |
-| 1920 | 1080 | 4  | 8.144ms  | 8.615ms  | 0.832ms   | 0.914ms   | 0.106ms  | 0.109ms  | 0.697ms  | 0.075ms |
-| 2560 | 1440 | 1  | 3.530ms  | 3.800ms  | 0.423ms   | 0.476ms   | 0.061ms  | 0.056ms  | 1.060ms  | 0.067ms |
-| 2560 | 1440 | 3  | 11.470ms | 11.611ms | 1.323ms   | 1.169ms   | 0.141ms  | 0.144ms  | 1.444ms  | 0.085ms |
-| 2560 | 1440 | 4  | 14.139ms | 15.273ms | 2.391ms   | 2.567ms   | 0.188ms  | 0.191ms  | 1.279ms  | 0.087ms |
-| 3840 | 2160 | 1  | 7.976ms  | 8.494ms  | 1.103ms   | 1.387ms   | 0.134ms  | 0.122ms  | 2.317ms  | 0.097ms |
-| 3840 | 2160 | 3  | 26.299ms | 25.824ms | 5.339ms   | 4.438ms   | 0.318ms  | 0.329ms  | 3.296ms  | 0.086ms |
-| 3840 | 2160 | 4  | 32.941ms | 34.718ms | 5.805ms   | 4.514ms   | 0.421ms  | 0.427ms  | 2.950ms  | 0.116ms |
-| 7680 | 4320 | 1  | 31.536ms | 34.100ms | 5.742ms   | 4.976ms   | 0.527ms  | 0.476ms  | 9.400ms  | 0.215ms |
-| 7680 | 4320 | 3  | 102.875ms| 102.419ms| 19.261ms  | 17.294ms  | 1.252ms  | 1.290ms  | 13.089ms | 0.223ms |
-| 7680 | 4320 | 4  | 133.081ms| 136.308ms| 23.398ms  | 18.445ms  | 1.670ms  | 1.688ms  | 11.529ms | 0.218ms |
+|       |     |    |    CPU   |    CPU   |CPU(OpenMP)|CPU(OpenMP)|   CUDA    |   CUDA  |   ROCm  |   ROCm  |
+|-------|-----|----|----------|----------|-----------|-----------|-----------|---------|---------|---------|
+|   W  |   H  |  C | hwc2chw  | chw2hwc  | hwc2chw   | chw2hwc   | hwc2chw   | chw2hwc | hwc2chw | chw2hwc |
+| 426  | 240  | 1  | 0.097ms  | 0.110ms  | 0.113ms   | 0.030ms   | 0.022ms   | 0.019ms | 0.059ms | 0.053ms |
+| 426  | 240  | 3  | 0.331ms  | 0.314ms  | 0.061ms   | 0.068ms   | 0.022ms   | 0.019ms | 0.062ms | 0.059ms |
+| 426  | 240  | 4  | 0.439ms  | 0.415ms  | 0.082ms   | 0.082ms   | 0.020ms   | 0.019ms | 0.062ms | 0.061ms |
+| 640  | 360  | 1  | 0.217ms  | 0.236ms  | 0.048ms   | 0.052ms   | 0.022ms   | 0.021ms | 0.062ms | 0.061ms |
+| 640  | 360  | 3  | 0.743ms  | 0.705ms  | 0.147ms   | 0.140ms   | 0.036ms   | 0.021ms | 0.060ms | 0.059ms |
+| 640  | 360  | 4  | 0.881ms  | 0.921ms  | 0.219ms   | 0.203ms   | 0.025ms   | 0.021ms | 0.057ms | 0.053ms |
+| 854  | 480  | 1  | 0.393ms  | 0.415ms  | 0.094ms   | 0.089ms   | 0.025ms   | 0.024ms | 0.063ms | 0.060ms |
+| 854  | 480  | 3  | 1.328ms  | 1.269ms  | 0.250ms   | 0.232ms   | 0.029ms   | 0.024ms | 0.052ms | 0.052ms |
+| 854  | 480  | 4  | 1.717ms  | 1.670ms  | 0.263ms   | 0.262ms   | 0.034ms   | 0.027ms | 0.054ms | 0.051ms |
+| 1280 | 720  | 1  | 0.873ms  | 0.937ms  | 0.130ms   | 0.180ms   | 0.053ms   | 0.040ms | 0.060ms | 0.052ms |
+| 1280 | 720  | 3  | 2.877ms  | 2.828ms  | 0.449ms   | 0.457ms   | 0.052ms   | 0.042ms | 0.061ms | 0.056ms |
+| 1280 | 720  | 4  | 3.558ms  | 3.848ms  | 0.719ms   | 0.616ms   | 0.054ms   | 0.045ms | 0.062ms | 0.056ms |
+| 1920 | 1080 | 1  | 1.949ms  | 2.136ms  | 0.374ms   | 0.342ms   | 0.081ms   | 0.067ms | 0.079ms | 0.060ms |
+| 1920 | 1080 | 3  | 6.587ms  | 6.469ms  | 1.000ms   | 0.672ms   | 0.087ms   | 0.074ms | 0.080ms | 0.064ms |
+| 1920 | 1080 | 4  | 8.144ms  | 8.615ms  | 0.832ms   | 0.914ms   | 0.103ms   | 0.080ms | 0.077ms | 0.057ms |
+| 2560 | 1440 | 1  | 3.530ms  | 3.800ms  | 0.423ms   | 0.476ms   | 0.114ms   | 0.116ms | 0.094ms | 0.074ms |
+| 2560 | 1440 | 3  | 11.470ms | 11.611ms | 1.323ms   | 1.169ms   | 0.142ms   | 0.127ms | 0.089ms | 0.070ms |
+| 2560 | 1440 | 4  | 14.139ms | 15.273ms | 2.391ms   | 2.567ms   | 0.154ms   | 0.136ms | 0.094ms | 0.075ms |
+| 3840 | 2160 | 1  | 7.976ms  | 8.494ms  | 1.103ms   | 1.387ms   | 0.234ms   | 0.227ms | 0.129ms | 0.097ms |
+| 3840 | 2160 | 3  | 26.299ms | 25.824ms | 5.339ms   | 4.438ms   | 0.307ms   | 0.253ms | 0.132ms | 0.096ms |
+| 3840 | 2160 | 4  | 32.941ms | 34.718ms | 5.805ms   | 4.514ms   | 0.323ms   | 0.272ms | 0.131ms | 0.097ms |
+| 7680 | 4320 | 1  | 31.536ms | 34.100ms | 5.742ms   | 4.976ms   | 0.836ms   | 0.741ms | 0.484ms | 0.214ms |
+| 7680 | 4320 | 3  | 102.875ms| 102.419ms| 19.261ms  | 17.294ms  | 1.057ms   | 0.890ms | 0.621ms | 0.222ms |
+| 7680 | 4320 | 4  | 133.081ms| 136.308ms| 23.398ms  | 18.445ms  | 1.144ms   | 1.013ms | 0.686ms | 0.220ms |
 
 ## Contact
 For any questions or suggestions, please open an issue or contact the me.
diff --git a/test/rocm_benchmark.cpp b/test/rocm_benchmark.cpp
@@ -42,15 +42,15 @@ int main() {
             //std::vector<uint8_t> out_uint8(pixel_size); // Inference output data(hwc)
 
             // 2. device memory
-            void* src_uint8 = 0;
-            void* src_float = 0;
-            hipHostMalloc(&src_uint8, pixel_size * sizeof(uint8_t), 0);
-            hipHostMalloc(&src_float, pixel_size * sizeof(float), 0);
+            hipDeviceptr_t src_uint8 = 0;
+            hipDeviceptr_t src_float = 0;
+            hipMalloc(&src_uint8, pixel_size * sizeof(uint8_t));
+            hipMalloc(&src_float, pixel_size * sizeof(float));
             
-            void* out_float = 0;
-            void* out_uint8 = 0;
-            hipHostMalloc(&out_float, pixel_size * sizeof(float), 0);
-            hipHostMalloc(&out_uint8, pixel_size * sizeof(uint8_t), 0);
+            hipDeviceptr_t out_float = 0;
+            hipDeviceptr_t out_uint8 = 0;
+            hipMalloc(&out_float, pixel_size * sizeof(float));
+            hipMalloc(&out_uint8, pixel_size * sizeof(uint8_t));
 
             auto startTime = std::chrono::high_resolution_clock::now();
             for (size_t i = 0; i < TEST_COUNT; ++i) {
@@ -79,10 +79,10 @@ int main() {
             auto chw2hwcDuration = std::chrono::duration_cast<std::chrono::microseconds>(endTime - startTime) / double(TEST_COUNT);
 
             // 2. device memory
-            hipHostFree(src_uint8);
-            hipHostFree(src_float);
-            hipHostFree(out_float);
-            hipHostFree(out_uint8);
+            hipFree(src_uint8);
+            hipFree(src_float);
+            hipFree(out_float);
+            hipFree(out_uint8);
 
             std::cout << width << ",\t" << height << ",\t" << channel << ",\t"
                 << std::fixed << std::setprecision(3)
