summaryrefslogtreecommitdiff
path: root/csrc/kernels.cu
diff options
context:
space:
mode:
authorTim Dettmers <tim.dettmers@gmail.com>2022-07-26 12:12:38 -0700
committerTim Dettmers <tim.dettmers@gmail.com>2022-07-26 12:12:38 -0700
commitcbb901ac51bd6c41e4243ffb936ef0e2f7ca8ada (patch)
treef02615b5588aa6ed94a51c1e66297595b802c0a1 /csrc/kernels.cu
parentc771b3a75a6ebbfbfc398a028a477246b0799cf0 (diff)
Boilerplate and test for extract_outliers.
Diffstat (limited to 'csrc/kernels.cu')
-rw-r--r--csrc/kernels.cu7
1 files changed, 7 insertions, 0 deletions
diff --git a/csrc/kernels.cu b/csrc/kernels.cu
index 1c3e723..78170d0 100644
--- a/csrc/kernels.cu
+++ b/csrc/kernels.cu
@@ -2592,10 +2592,17 @@ __global__ void kspmm_coo_very_sparse_naive(int *max_count, int *max_idx, int *o
}
}
+template <int FORMAT> __global__ void kExtractOutliers(char *A, int *idx, char *out, int rowsA, int colsA, int tiledRowsA, int tiledColsA)
+{
+}
+
//==============================================================
// TEMPLATE DEFINITIONS
//==============================================================
+template __global__ void kExtractOutliers<COL_TURING>(char *A, int *idx, char *out, int rowsA, int colsA, int tiledRowsA, int tiledColsA);
+template __global__ void kExtractOutliers<COL_AMPERE>(char *A, int *idx, char *out, int rowsA, int colsA, int tiledRowsA, int tiledColsA);
+
template __global__ void kspmm_coo_very_sparse_naive<half, 8, 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, int rowsA, int rowsB, int colsB);
template __global__ void kspmm_coo_very_sparse_naive<half, 16, 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, int rowsA, int rowsB, int colsB);
template __global__ void kspmm_coo_very_sparse_naive<half, 32, 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, int rowsA, int rowsB, int colsB);