summaryrefslogtreecommitdiff
path: root/csrc/kernels.cuh
diff options
context:
space:
mode:
authorTim Dettmers <tim.dettmers@gmail.com>2022-08-16 10:56:17 -0700
committerTim Dettmers <tim.dettmers@gmail.com>2022-08-16 10:56:17 -0700
commit1ed2fa2f218d8dac401f3315420ffec92014c124 (patch)
tree57863d4d1024689100c1b43caccc1d8739c58d99 /csrc/kernels.cuh
parent26efb154c8d77b4ede2cfc0dbd2381dd385f33e7 (diff)
Removed storage() from get_ptr; added boilerplate for bias dequant_mm.
Diffstat (limited to 'csrc/kernels.cuh')
-rw-r--r--csrc/kernels.cuh2
1 files changed, 1 insertions, 1 deletions
diff --git a/csrc/kernels.cuh b/csrc/kernels.cuh
index 2447494..bdf61b2 100644
--- a/csrc/kernels.cuh
+++ b/csrc/kernels.cuh
@@ -111,7 +111,7 @@ template <typename T, int SPMM_ITEMS, int BITS> __global__ void kspmm_coo_very_s
template <int ITEMS_PER_THREAD, int SUBTILE_ROWS, int THREADS>__global__ void kdequant_mm_int32_fp16(
int *__restrict__ const A, float *__restrict__ const rowStats, float *__restrict__ const colStats,
- half *out, float* newRowStats, float* newcolStats, const int numRows, const int numCols, const int tileCols, const int n);
+ half *out, float* newRowStats, float* newcolStats, half * __restrict__ const bias, const int numRows, const int numCols, const int tileCols, const int n);
template<typename T, int THREADS, int ITEMS_PER_THREAD, int TILE_ROWS, int TILE_COLS, int SPARSE_DECOMP> __global__ void kgetColRowStats(T * __restrict__ A, float *rowStats, float *colStats, int * nnz_count_row, float nnz_threshold, int rows, int cols, int tiledRows, int tiledCols);
template <int THREADS, int ITEMS_PER_THREAD, int TILE_ROWS, int TILE_COLS, int SPARSE_DECOMP> __global__ void kDoubleRowColQuant(half *__restrict__ const A, float *__restrict__ const rowStats, float * __restrict__ const colStats, char *out_col_normed, char *out_row_normed, int *rowidx, int *colidx, half *val, int * __restrict__ nnz_block_ptr, float threshold, int rows, int cols, int tiledCols);