Merge pull request #227 from sampath1117/sr/histogram_build_fix

Histogram build fix
r-abishek · Jan 29, 2024 · 363e305 · 363e305
2 parents 2d7adde + 80ea779
commit 363e305
Showing 1 changed file with 12 additions and 12 deletions.
diff --git a/src/modules/hip/kernel/histogram.cpp b/src/modules/hip/kernel/histogram.cpp
@@ -39,9 +39,9 @@ void partial_histogram_pln( unsigned char *input,
         unsigned char pixelR = input[pixId];
         unsigned char pixelG = input[pixId + width * height];
         unsigned char pixelB = input[pixId + 2 * width * height];
-        atomicInc(&tmp_histogram[pixelR], 1);
-        atomicInc(&tmp_histogram[pixelG], 1);
-        atomicInc(&tmp_histogram[pixelB], 1);
+        atomicAdd(&tmp_histogram[pixelR], 1);
+        atomicAdd(&tmp_histogram[pixelG], 1);
+        atomicAdd(&tmp_histogram[pixelB], 1);
     }
     __syncthreads();
     if (local_size >= (256 ))
@@ -98,9 +98,9 @@ void partial_histogram_pkd( unsigned char *input,
         unsigned char pixelR = input[pixId];
         unsigned char pixelG = input[pixId + 1];
         unsigned char pixelB = input[pixId + 2];
-        atomicInc(&tmp_histogram[pixelR], 1);
-        atomicInc(&tmp_histogram[pixelG], 1);
-        atomicInc(&tmp_histogram[pixelB], 1);
+        atomicAdd(&tmp_histogram[pixelR], 1);
+        atomicAdd(&tmp_histogram[pixelG], 1);
+        atomicAdd(&tmp_histogram[pixelB], 1);
     }
     __syncthreads();
     if (local_size >= (256 ))
@@ -166,9 +166,9 @@ void partial_histogram_batch( unsigned char* input,
         unsigned char pixelR = input[pixId];
         unsigned char pixelG = input[pixId + inc[id_z]];
         unsigned char pixelB = input[pixId + 2 * inc[id_z]];
-        atomicInc(&tmp_histogram[pixelR], 1);
-        atomicInc(&tmp_histogram[pixelG], 1);
-        atomicInc(&tmp_histogram[pixelB], 1);
+        atomicAdd(&tmp_histogram[pixelR], 1);
+        atomicAdd(&tmp_histogram[pixelG], 1);
+        atomicAdd(&tmp_histogram[pixelB], 1);
     }
     __syncthreads();
     if (local_size >= (256 ))
@@ -234,9 +234,9 @@ void partial_histogram_semibatch( unsigned char* input,
         unsigned char pixelR = input[pixId];
         unsigned char pixelG = input[pixId + inc];
         unsigned char pixelB = input[pixId + 2 * inc];
-       atomicInc(&tmp_histogram[pixelR], 1);
-        atomicInc(&tmp_histogram[pixelG], 1);
-        atomicInc(&tmp_histogram[pixelB], 1);
+       atomicAdd(&tmp_histogram[pixelR], 1);
+        atomicAdd(&tmp_histogram[pixelG], 1);
+        atomicAdd(&tmp_histogram[pixelB], 1);
     }
     __syncthreads();
     if (local_size >= (256 ))