[SYCL-MLIR] Manual optimizations (Scalar replacement)

Signed-off-by: Tsang, Whitney <[email protected]>
whitneywhtsang · Feb 18, 2023 · 8121a14 · 8121a14
1 parent ef03944
commit 8121a14
Show file tree

Hide file tree

Showing 11 changed files with 59 additions and 23 deletions.
diff --git a/polybench/2mm.cpp b/polybench/2mm.cpp
@@ -100,9 +100,11 @@ class Polybench_2mm {
 				const auto i = item[0];
 				const auto j = item[1];
 
+        DATA_TYPE C_reduction = C[item];
 				for(size_t k = 0; k < size_; k++) {
-					C[item] += A[{i, k}] * B[{k, j}];
+					C_reduction += A[{i, k}] * B[{k, j}];
 				}
+				C[item] = C_reduction;
 			});
 		}));
 
@@ -116,9 +118,11 @@ class Polybench_2mm {
 				const auto j = item[1];
 
 				E[item] = 0;
+        DATA_TYPE E_reduction = E[item];
 				for(size_t k = 0; k < size_; k++) {
-					E[item] += C[{i, k}] * D[{k, j}];
+					E_reduction += C[{i, k}] * D[{k, j}];
 				}
+				E[item] = E_reduction;
 			});
 		}));
 	}

diff --git a/polybench/3mm.cpp b/polybench/3mm.cpp
@@ -120,9 +120,11 @@ class Polybench_3mm {
 				const auto i = item[0];
 				const auto j = item[1];
 
+        DATA_TYPE E_reduction = E[item];
 				for(size_t k = 0; k < size_; k++) {
-					E[item] += A[{i, k}] * B[{k, j}];
+					E_reduction += A[{i, k}] * B[{k, j}];
 				}
+				E[item] = E_reduction;
 			});
 		}));
 
@@ -135,9 +137,11 @@ class Polybench_3mm {
 				const auto i = item[0];
 				const auto j = item[1];
 
+        DATA_TYPE F_reduction = F[item];
 				for(size_t k = 0; k < size_; k++) {
-					F[item] += C[{i, k}] * D[{k, j}];
+					F_reduction += C[{i, k}] * D[{k, j}];
 				}
+				F[item] = F_reduction;
 			});
 		}));
 
@@ -150,9 +154,11 @@ class Polybench_3mm {
 				const auto i = item[0];
 				const auto j = item[1];
 
+        DATA_TYPE G_reduction = G[item];
 				for(size_t k = 0; k < size_; k++) {
-					G[item] += E[{i, k}] * F[{k, j}];
+					G_reduction += E[{i, k}] * F[{k, j}];
 				}
+				G[item] = G_reduction;
 			});
 		}));
 	}

diff --git a/polybench/atax.cpp b/polybench/atax.cpp
@@ -73,9 +73,11 @@ class Polybench_Atax {
 			cgh.parallel_for<Atax1>(tmp_buffer.get_range(), [=, size_ = size](item<1> item) {
 				const auto i = item[0];
 
+        DATA_TYPE tmp_reduction = tmp[item];
 				for(size_t j = 0; j < size_; j++) {
-					tmp[item] += A[{i, j}] * x[j];
+					tmp_reduction += A[{i, j}] * x[j];
 				}
+				tmp[item] = tmp_reduction;
 			});
 		}));
 
@@ -87,9 +89,11 @@ class Polybench_Atax {
 			cgh.parallel_for<Atax2>(y_buffer.get_range(), [=, size_ = size](item<1> item) {
 				const auto j = item[0];
 
+        DATA_TYPE y_reduction = y[item];
 				for(size_t i = 0; i < size_; i++) {
-					y[item] += A[{i, j}] * tmp[i];
+					y_reduction += A[{i, j}] * tmp[i];
 				}
+				y[item] = y_reduction;
 			});
 		}));
 	}

diff --git a/polybench/bicg.cpp b/polybench/bicg.cpp
@@ -77,9 +77,11 @@ class Polybench_Bicg {
 			cgh.parallel_for<Bicg1>(s_buffer.get_range(), [=, size_ = size](item<1> item) {
 				const auto j = item[0];
 
+        DATA_TYPE s_reduction = s[item];
 				for(size_t i = 0; i < size_; i++) {
-					s[item] += A[{i, j}] * r[i];
+					s_reduction += A[{i, j}] * r[i];
 				}
+				s[item] = s_reduction;
 			});
 		}));
 
@@ -91,9 +93,11 @@ class Polybench_Bicg {
 			cgh.parallel_for<Bicg2>(q_buffer.get_range(), [=, size_ = size](item<1> item) {
 				const auto i = item[0];
 
+        DATA_TYPE q_reduction = q[item];
 				for(size_t j = 0; j < size_; j++) {
-					q[item] += A[{i, j}] * p[j];
+					q_reduction += A[{i, j}] * p[j];
 				}
+				q[item] = q_reduction;
 			});
 		}));
 	}

diff --git a/polybench/covariance.cpp b/polybench/covariance.cpp
@@ -86,10 +86,12 @@ class Polybench_Covariance {
 			cgh.parallel_for<CovarianceMean>(range<1>(size), id<1>(1), [=, N_ = size](item<1> item) {
 				const auto j = item[0];
 
-				mean[item] = 0;
+        mean[item] = 0;
+        DATA_TYPE mean_reduction = mean[item];
 				for(size_t i = 1; i <= N_; i++) {
-					mean[item] += data[{i, j}];
+					mean_reduction += data[{i, j}];
 				}
+				mean[item] = mean_reduction;
 				mean[item] /= float_n;
 			});
 		}));
@@ -115,10 +117,12 @@ class Polybench_Covariance {
 				symmat[{j1, j1}] = 1.0;
 
 				for(size_t j2 = j1; j2 <= M_; j2++) {
-					symmat[{j1, j2}] = 0.0;
+          symmat[{j1, j2}] = 0.0;
+					DATA_TYPE symmat_reduction = symmat[{j1, j2}];
 					for(size_t i = 1; i <= N_; i++) {
-						symmat[{j1, j2}] += data[{i, j1}] * data[{i, j2}];
+						symmat_reduction += data[{i, j1}] * data[{i, j2}];
 					}
+					symmat[{j1, j2}] = symmat_reduction;
 
 					symmat2[{j2, j1}] = symmat[{j1, j2}];
 				}

diff --git a/polybench/gemm.cpp b/polybench/gemm.cpp
@@ -84,10 +84,11 @@ class Polybench_Gemm {
 				const auto j = item[1];
 
 				C[item] *= BETA;
-
+        DATA_TYPE C_reduction = C[item];
 				for(size_t k = 0; k < NK_; k++) {
-					C[item] += ALPHA * A[{i, k}] * B[{k, j}];
+					C_reduction += ALPHA * A[{i, k}] * B[{k, j}];
 				}
+				C[item] = C_reduction;
 			});
 		}));
 	}

diff --git a/polybench/gesummv.cpp b/polybench/gesummv.cpp
@@ -87,10 +87,14 @@ class Polybench_Gesummv {
 			cgh.parallel_for<Gesummv>(y.get_range(), [=, N_ = size](item<1> item) {
 				const auto i = item[0];
 
+        DATA_TYPE tmp_reduction = tmp[item];
+				DATA_TYPE y_reduction = y[item];
 				for(size_t j = 0; j < N_; j++) {
-					tmp[item] += A[{i, j}] * x[j];
-					y[item] += B[{i, j}] * x[j];
+					tmp_reduction += A[{i, j}] * x[j];
+					y_reduction += B[{i, j}] * x[j];
 				}
+				tmp[item] = tmp_reduction;
+				y[item] = y_reduction;
 
 				y[item] = ALPHA * tmp[item] + BETA * y[item];
 			});

diff --git a/polybench/gramschmidt.cpp b/polybench/gramschmidt.cpp
@@ -104,10 +104,11 @@ class Polybench_Gramschmidt {
 
 					if(j <= k || j >= N_) return;
 
-					R[item] = 0;
+					DATA_TYPE R_reduction = 0;
 					for(size_t i = 0; i < M_; i++) {
-						R[item] += Q[{i, k}] * A[{i, j}];
+						R_reduction += Q[{i, k}] * A[{i, j}];
 					}
+					R[item] = R_reduction;
 
 					for(size_t i = 0; i < M_; i++) {
 						A[{i, j}] -= Q[{i, k}] * R[item];

diff --git a/polybench/mvt.cpp b/polybench/mvt.cpp
@@ -75,9 +75,11 @@ class Polybench_Mvt {
 			cgh.parallel_for<Mvt1>(x1_buffer.get_range(), [=, N_ = size](item<1> item) {
 				const auto i = item[0];
 
+        DATA_TYPE x1_reduction = x1[i];
 				for(size_t j = 0; j < N_; j++) {
-					x1[i] += a[{i, j}] * y1[j];
+					x1_reduction += a[{i, j}] * y1[j];
 				}
+				x1[i] = x1_reduction;
 			});
 		}));
 
@@ -89,9 +91,11 @@ class Polybench_Mvt {
 			cgh.parallel_for<Mvt2>(x1_buffer.get_range(), [=, N_ = size](item<1> item) {
 				const auto k = item[0];
 
+        DATA_TYPE x2_reduction = x2[k];
 				for(size_t l = 0; l < N_; l++) {
-					x2[k] += a[{k, l}] * y2[l];
+					x2_reduction += a[{k, l}] * y2[l];
 				}
+				x2[k] = x2_reduction;
 			});
 		}));
 	}

diff --git a/polybench/syr2k.cpp b/polybench/syr2k.cpp
@@ -81,9 +81,11 @@ class Polybench_Syr2k {
 
 				C[item] *= BETA;
 
+        DATA_TYPE C_reduction = C[item];
 				for(size_t k = 0; k < M_; k++) {
-					C[item] += ALPHA * A[{i, k}] * B[{j, k}] + ALPHA * B[{i, k}] * A[{j, k}];
+					C_reduction += ALPHA * A[{i, k}] * B[{j, k}] + ALPHA * B[{i, k}] * A[{j, k}];
 				}
+				C[item] = C_reduction;
 			});
 		}));
 	}

diff --git a/polybench/syrk.cpp b/polybench/syrk.cpp
@@ -77,9 +77,11 @@ class Polybench_Syrk {
 
 				C[item] *= beta;
 
+        DATA_TYPE C_reduction = C[item];
 				for(size_t k = 0; k < M_; k++) {
-					C[item] += alpha * A[{i, k}] * A[{j, k}];
+					C_reduction += alpha * A[{i, k}] * A[{j, k}];
 				}
+				C[item] = C_reduction;
 			});
 		}));
 	}