From cbb901ac51bd6c41e4243ffb936ef0e2f7ca8ada Mon Sep 17 00:00:00 2001
From: Tim Dettmers <tim.dettmers@gmail.com>
Date: Tue, 26 Jul 2022 12:12:38 -0700
Subject: Boilerplate and test for extract_outliers.

---
 csrc/ops.cu | 27 +++++++++++++++++++++++++++
 1 file changed, 27 insertions(+)

(limited to 'csrc/ops.cu')

diff --git a/csrc/ops.cu b/csrc/ops.cu
index 8946015..fe2d7fe 100644
--- a/csrc/ops.cu
+++ b/csrc/ops.cu
@@ -578,10 +578,37 @@ template <typename T, int BITS> void spmm_coo_very_sparse_naive(int *max_count,
   CUDA_CHECK_RETURN(cudaPeekAtLastError());
 }
 
+
+template <int FORMAT> void extractOutliers(char * A, int *idx, char *out, int idx_size, int rows, int cols)
+{
+  int threads = 256;
+  // we load 128 column values per warp
+  int tiledCols = tiledCols = fill_up_to_nearest_multiple(cols, 32);
+  int tiledRows = 0;
+
+	int elements = idx_size*cols; // matrix A is transposed, so we extract columns
+	int num_blocks = (elements+threads-1)/threads;
+
+  if(FORMAT == COL_TURING)
+  {
+      tiledRows = fill_up_to_nearest_multiple(rows, 8);
+  }
+  else if(FORMAT == COL_AMPERE)
+  {
+      tiledRows = fill_up_to_nearest_multiple(rows, 32);
+	}
+
+  kExtractOutliers<FORMAT><<<num_blocks, threads>>>(A, idx, out, rows, cols, tiledRows, tiledCols);
+  CUDA_CHECK_RETURN(cudaPeekAtLastError());
+}
+
 //==============================================================
 //                   TEMPLATE DEFINITIONS
 //==============================================================
 
+template void extractOutliers<COL_TURING>(char * A, int *idx, char *out, int idx_size, int rows, int cols);
+template void extractOutliers<COL_AMPERE>(char * A, int *idx, char *out, int idx_size, int rows, int cols);
+
 template void spmm_coo_very_sparse_naive<half, 16>(int *max_count, int *max_idx, int *offset_rowidx, int *rowidx, int *colidx, half *values, half *B, half *out, float *dequant_stats, int nnz_rows, int nnz, int rowsA, int rowsB, int colsB);
 template void spmm_coo_very_sparse_naive<signed char, 8>(int *max_count, int *max_idx, int *offset_rowidx, int *rowidx, int *colidx, half *values, signed char *B, half *out, float *dequant_stats, int nnz_rows, int nnz, int rowsA, int rowsB, int colsB);
 
-- 
cgit v1.2.3


From bcab99ec877ba063543bd7c03ba1cdd1b06e8078 Mon Sep 17 00:00:00 2001
From: Tim Dettmers <tim.dettmers@gmail.com>
Date: Tue, 26 Jul 2022 17:39:30 -0700
Subject: Working outlier extraction for Turing.

---
 csrc/ops.cu | 5 ++---
 1 file changed, 2 insertions(+), 3 deletions(-)

(limited to 'csrc/ops.cu')

diff --git a/csrc/ops.cu b/csrc/ops.cu
index fe2d7fe..e6227ae 100644
--- a/csrc/ops.cu
+++ b/csrc/ops.cu
@@ -586,8 +586,7 @@ template <int FORMAT> void extractOutliers(char * A, int *idx, char *out, int id
   int tiledCols = tiledCols = fill_up_to_nearest_multiple(cols, 32);
   int tiledRows = 0;
 
-	int elements = idx_size*cols; // matrix A is transposed, so we extract columns
-	int num_blocks = (elements+threads-1)/threads;
+	int num_blocks = idx_size;
 
   if(FORMAT == COL_TURING)
   {
@@ -598,7 +597,7 @@ template <int FORMAT> void extractOutliers(char * A, int *idx, char *out, int id
       tiledRows = fill_up_to_nearest_multiple(rows, 32);
 	}
 
-  kExtractOutliers<FORMAT><<<num_blocks, threads>>>(A, idx, out, rows, cols, tiledRows, tiledCols);
+  kExtractOutliers<FORMAT><<<num_blocks, threads>>>(A, idx, out, idx_size, rows, cols, tiledRows, tiledCols);
   CUDA_CHECK_RETURN(cudaPeekAtLastError());
 }
 
-- 
cgit v1.2.3