1. Fix compile warning.

2. Modify kernal function.
Signed-off-by: lijian <lijian6@sugon.com>

CMakeLists.txt

@@ -2,7 +2,7 @@ cmake_minimum_required(VERSION 3.16)
 project(Video_MIGraphX)
 
 set(CMAKE_CXX_COMPILER hipcc)
-set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++17 -pthread -lz")
+set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++17 -pthread")
 set(CMAKE_BUILD_TYPE release)
 
 set(INCLUDE_PATH    /usr/local/include

src/Inference/DetectorRetinaFace.cpp

@@ -223,9 +223,9 @@ __global__ void convert_yuv420p_to_rgb_and_normalization_retinaface(unsigned cha
     int g = yValue - 0.698001 * (vValue - 128) - 0.337633 * (uValue - 128);
     int b = yValue + 1.732446 * (uValue - 128);
 
-    outImage[index + 0] = (float)r;
-    outImage[index + 1] = (float)g;
-    outImage[index + 2] = (float)b;
+    outImage[x + width * height * 0] = (float)r;
+    outImage[x + width * height * 1] = (float)g;
+    outImage[x + width * height * 2] = (float)b;
 }
 __global__ void convert_rgba_to_rgb_and_normalization_retinaface(unsigned char* srcImage, float* outImage, int width, int height)
 {

src/Inference/DetectorSSD.cpp

@@ -224,9 +224,9 @@ __global__ void convert_yuv420p_to_rgb_and_normalization_ssd(unsigned char* srcI
     int g = yValue - 0.698001 * (vValue - 128) - 0.337633 * (uValue - 128);
     int b = yValue + 1.732446 * (uValue - 128);
 
-    outImage[index + 0] = (float)r;
-    outImage[index + 1] = (float)g;
-    outImage[index + 2] = (float)b;
+    outImage[x + width * height * 0] = (float)r;
+    outImage[x + width * height * 1] = (float)g;
+    outImage[x + width * height * 2] = (float)b;
 }
 __global__ void convert_rgba_to_rgb_and_normalization_ssd(unsigned char* srcImage, float* outImage, int width, int height)
 {

src/Inference/DetectorYOLOV3.cpp

@@ -160,9 +160,9 @@ __global__ void convert_yuv420p_to_rgb_and_normalization_yolov3(unsigned char* s
     int g = yValue - 0.698001 * (vValue - 128) - 0.337633 * (uValue - 128);
     int b = yValue + 1.732446 * (uValue - 128);
 
-    outImage[index + 0] = (float)r;
-    outImage[index + 1] = (float)g;
-    outImage[index + 2] = (float)b;
+    outImage[x + width * height * 0] = (float)r;
+    outImage[x + width * height * 1] = (float)g;
+    outImage[x + width * height * 2] = (float)b;
 }
 __global__ void convert_rgba_to_rgb_and_normalization_yolov3(unsigned char* srcImage, float* outImage, int width, int height)
 {

src/Inference/DetectorYOLOV5.cpp

@@ -160,9 +160,9 @@ __global__ void convert_yuv420p_to_rgb_and_normalization_yolov5(unsigned char* s
     int g = yValue - 0.698001 * (vValue - 128) - 0.337633 * (uValue - 128);
     int b = yValue + 1.732446 * (uValue - 128);
 
-    outImage[index + 0] = (float)r;
-    outImage[index + 1] = (float)g;
-    outImage[index + 2] = (float)b;
+    outImage[x + width * height * 0] = (float)r;
+    outImage[x + width * height * 1] = (float)g;
+    outImage[x + width * height * 2] = (float)b;
 }
 __global__ void convert_rgba_to_rgb_and_normalization_yolov5(unsigned char* srcImage, float* outImage, int width, int height)
 {

src/Inference/DetectorYOLOV7.cpp

@@ -159,9 +159,9 @@ __global__ void convert_yuv420p_to_rgb_and_normalization_yolov7(unsigned char* s
     int g = yValue - 0.698001 * (vValue - 128) - 0.337633 * (uValue - 128);
     int b = yValue + 1.732446 * (uValue - 128);
 
-    outImage[index + 0] = (float)r;
-    outImage[index + 1] = (float)g;
-    outImage[index + 2] = (float)b;
+    outImage[x + width * height * 0] = (float)r;
+    outImage[x + width * height * 1] = (float)g;
+    outImage[x + width * height * 2] = (float)b;
 }
 __global__ void convert_rgba_to_rgb_and_normalization_yolov7(unsigned char* srcImage, float* outImage, int width, int height)
 {

src/RetinaFace.cpp

@@ -22,7 +22,7 @@ static void DecoderThreadFunc(Queue* queue)
     #ifndef DMA
     initParamOfDecoderRetinaFace.src_filename = "../Resource/Images/Mean.mp4";
     if( que->device == _HW) {
-        initParamOfDecoderRetinaFace.str_devid[4] = {0};
+        initParamOfDecoderRetinaFace.str_devid[0] = {0};
         initParamOfDecoderRetinaFace.xcoder_params = "out=hw";
         initParamOfDecoderRetinaFace.dec_name = "h264_ni_quadra_dec";
         initParamOfDecoderRetinaFace.filters_descr = "ni_quadra_scale=640:640:format=bgrp,hwdownload,format=bgrp";
@@ -33,7 +33,7 @@ static void DecoderThreadFunc(Queue* queue)
     }
     #else
     if( que->device == _HW_DMA) {
-        initParamOfDecoderRetinaFace.str_devid[4] = {0};
+        initParamOfDecoderRetinaFace.str_devid[0] = {0};
         initParamOfDecoderRetinaFace.xcoder_params = "out=hw";
         initParamOfDecoderRetinaFace.dec_name = "h264_ni_quadra_dec";
         initParamOfDecoderRetinaFace.filters_descr = "ni_quadra_scale=640:640:format=rgba:is_p2p=1";

src/SSD.cpp

@@ -22,7 +22,7 @@ static void DecoderThreadFunc(Queue* queue)
     #ifndef DMA
     initParamOfDecoderSSD.src_filename = "../Resource/Images/Mean.mp4";
     if( que->device == _HW) {
-        initParamOfDecoderSSD.str_devid[4] = {0};
+        initParamOfDecoderSSD.str_devid[0] = {0};
         initParamOfDecoderSSD.xcoder_params = "out=hw";
         initParamOfDecoderSSD.dec_name = "h264_ni_quadra_dec";
         initParamOfDecoderSSD.filters_descr = "ni_quadra_scale=640:480:format=bgrp,hwdownload,format=bgrp";
@@ -33,7 +33,7 @@ static void DecoderThreadFunc(Queue* queue)
     }
     #else
     if( que->device == _HW_DMA) {
-        initParamOfDecoderSSD.str_devid[4] = {0};
+        initParamOfDecoderSSD.str_devid[0] = {0};
         initParamOfDecoderSSD.xcoder_params = "out=hw";
         initParamOfDecoderSSD.dec_name = "h264_ni_quadra_dec";
         initParamOfDecoderSSD.filters_descr = "ni_quadra_scale=640:480:format=rgba:is_p2p=1";

src/YOLOV5.cpp

@@ -22,7 +22,7 @@ static void DecoderThreadFunc(Queue* queue)
     #ifndef DMA
     initParamOfDecoderYOLOV5.src_filename = "../Resource/Images/Mean.mp4";
     if( que->device == _HW) {
-        initParamOfDecoderYOLOV5.str_devid[4] = {0};
+        initParamOfDecoderYOLOV5.str_devid[0] = {0};
         initParamOfDecoderYOLOV5.xcoder_params = "out=hw";
         initParamOfDecoderYOLOV5.dec_name = "h264_ni_quadra_dec";
         initParamOfDecoderYOLOV5.filters_descr = "ni_quadra_scale=608:608:format=bgrp,hwdownload,format=bgrp";
@@ -33,7 +33,7 @@ static void DecoderThreadFunc(Queue* queue)
     }
     #else
     if( que->device == _HW_DMA) {
-        initParamOfDecoderYOLOV5.str_devid[4] = {0};
+        initParamOfDecoderYOLOV5.str_devid[0] = {0};
         initParamOfDecoderYOLOV5.xcoder_params = "out=hw";
         initParamOfDecoderYOLOV5.dec_name = "h264_ni_quadra_dec";
         initParamOfDecoderYOLOV5.filters_descr = "ni_quadra_scale=608:608:format=rgba:is_p2p=1";

src/YOLOV7.cpp

@@ -23,7 +23,7 @@ static void DecoderThreadFunc(Queue* queue)
     initParamOfDecoderYOLOV7.src_filename = "../Resource/Images/Mean.mp4";
     if( que->device == _HW)
     {
-	initParamOfDecoderYOLOV7.str_devid[4] = {0};
+	initParamOfDecoderYOLOV7.str_devid[0] = {0};
 	initParamOfDecoderYOLOV7.xcoder_params = "out=hw";
 	initParamOfDecoderYOLOV7.dec_name = "h264_ni_quadra_dec";
 	initParamOfDecoderYOLOV7.filters_descr = "ni_quadra_scale=640:640:format=bgrp,hwdownload,format=bgrp";
@@ -34,7 +34,7 @@ static void DecoderThreadFunc(Queue* queue)
     }
     #else
     if( que->device == _HW_DMA) {
-	initParamOfDecoderYOLOV7.str_devid[4] = {0};
+	initParamOfDecoderYOLOV7.str_devid[0] = {0};
 	initParamOfDecoderYOLOV7.xcoder_params = "out=hw";
 	initParamOfDecoderYOLOV7.dec_name = "h264_ni_quadra_dec";
 	initParamOfDecoderYOLOV7.filters_descr = "ni_quadra_scale=640:640:format=rgba:is_p2p=1";