kmeans f64 emulation

Author: Hardcode84

dpbench/benchmarks/kmeans/kmeans_numba_mlir_k.py

@@ -9,7 +9,7 @@
 atomic_add = nb.atomic.add
 
 
-@nb.kernel
+@nb.kernel(gpu_fp64_truncate="auto")
 def groupByCluster(arrayP, arrayPcluster, arrayC, num_points, num_centroids):
     idx = nb.get_global_id(0)
     # if idx < num_points: # why it was removed??
@@ -23,15 +23,15 @@ def groupByCluster(arrayP, arrayPcluster, arrayC, num_points, num_centroids):
             arrayPcluster[idx] = i
 
 
-@nb.kernel
+@nb.kernel(gpu_fp64_truncate="auto")
 def calCentroidsSum1(arrayCsum, arrayCnumpoint):
     i = nb.get_global_id(0)
     arrayCsum[i, 0] = 0
     arrayCsum[i, 1] = 0
     arrayCnumpoint[i] = 0
 
 
-@nb.kernel
+@nb.kernel(gpu_fp64_truncate="auto")
 def calCentroidsSum2(arrayP, arrayPcluster, arrayCsum, arrayCnumpoint):
     i = nb.get_global_id(0)
     ci = arrayPcluster[i]
@@ -40,14 +40,14 @@ def calCentroidsSum2(arrayP, arrayPcluster, arrayCsum, arrayCnumpoint):
     atomic_add(arrayCnumpoint, ci, 1)
 
 
-@nb.kernel
+@nb.kernel(gpu_fp64_truncate="auto")
 def updateCentroids(arrayC, arrayCsum, arrayCnumpoint, num_centroids):
     i = nb.get_global_id(0)
     arrayC[i, 0] = arrayCsum[i, 0] / arrayCnumpoint[i]
     arrayC[i, 1] = arrayCsum[i, 1] / arrayCnumpoint[i]
 
 
-@nb.kernel
+@nb.kernel(gpu_fp64_truncate="auto")
 def copy_arrayC(arrayC, arrayP):
     i = nb.get_global_id(0)
     arrayC[i, 0] = arrayP[i, 0]

dpbench/benchmarks/kmeans/kmeans_numba_mlir_p.py

@@ -11,7 +11,7 @@
 
 
 # determine the euclidean distance from the cluster center to each point
-@nb.njit
+@nb.njit(parallel=True, fastmath=True, gpu_fp64_truncate="auto")
 def groupByCluster(arrayP, arrayPcluster, arrayC, num_points, num_centroids):
     # parallel for loop
     for i0 in numba.prange(num_points):
@@ -27,7 +27,7 @@ def groupByCluster(arrayP, arrayPcluster, arrayC, num_points, num_centroids):
 
 
 # assign points to cluster
-@nb.njit
+@nb.njit(parallel=True, fastmath=True, gpu_fp64_truncate="auto")
 def calCentroidsSum(
     arrayP, arrayPcluster, arrayCsum, arrayCnumpoint, num_points, num_centroids
 ):
@@ -38,7 +38,7 @@ def calCentroidsSum(
         arrayCnumpoint[i] = 0
 
 
-@nbk.kernel
+@nbk.kernel(gpu_fp64_truncate="auto")
 def calCentroidsSum2(arrayP, arrayPcluster, arrayCsum, arrayCnumpoint):
     i = nbk.get_global_id(0)
     ci = arrayPcluster[i]
@@ -48,14 +48,14 @@ def calCentroidsSum2(arrayP, arrayPcluster, arrayCsum, arrayCnumpoint):
 
 
 # update the centriods array after computation
-@nb.njit
+@nb.njit(parallel=True, fastmath=True, gpu_fp64_truncate="auto")
 def updateCentroids(arrayC, arrayCsum, arrayCnumpoint, num_centroids):
     for i in numba.prange(num_centroids):
         arrayC[i, 0] = arrayCsum[i, 0] / arrayCnumpoint[i]
         arrayC[i, 1] = arrayCsum[i, 1] / arrayCnumpoint[i]
 
 
-@nb.njit
+@nb.njit(parallel=True, fastmath=True, gpu_fp64_truncate="auto")
 def copy_arrayC(arrayC, arrayP, num_centroids):
     for i in numba.prange(num_centroids):
         arrayC[i, 0] = arrayP[i, 0]