Lines Matching refs:sliceheight
75 static __global__ void matmult_seqsell_basic_kernel(PetscInt nrows, PetscInt sliceheight, const Pet… in matmult_seqsell_basic_kernel() argument
82 slice_id = row / sliceheight; in matmult_seqsell_basic_kernel()
83 row_in_slice = row % sliceheight; in matmult_seqsell_basic_kernel()
85 …for (i = sliidx[slice_id] + row_in_slice; i < sliidx[slice_id + 1]; i += sliceheight) sum += aval[… in matmult_seqsell_basic_kernel()
90 static __global__ void matmultadd_seqsell_basic_kernel(PetscInt nrows, PetscInt sliceheight, const … in matmultadd_seqsell_basic_kernel() argument
97 slice_id = row / sliceheight; in matmultadd_seqsell_basic_kernel()
98 row_in_slice = row % sliceheight; in matmultadd_seqsell_basic_kernel()
100 …for (i = sliidx[slice_id] + row_in_slice; i < sliidx[slice_id + 1]; i += sliceheight) sum += aval[… in matmultadd_seqsell_basic_kernel()
108 __global__ void matmult_seqsell_tiled_kernel9(PetscInt nrows, PetscInt sliceheight, const PetscInt … in matmult_seqsell_tiled_kernel9() argument
118 row = slice_id * sliceheight + threadIdx.x % sliceheight; in matmult_seqsell_tiled_kernel9()
123 …for (int offset = 16; offset >= sliceheight; offset /= 2) t += __shfl_down_sync(0xffffffff, t, off… in matmult_seqsell_tiled_kernel9()
125 if (threadIdx.x < sliceheight) shared[threadIdx.x][threadIdx.y] = t; in matmult_seqsell_tiled_kernel9()
127 if (tidy < sliceheight) t = shared[tidy][tidx]; in matmult_seqsell_tiled_kernel9()
130 if (tidx == 0 && tidy < sliceheight) shared[0][tidy] = t; in matmult_seqsell_tiled_kernel9()
132 if (row < nrows && threadIdx.y == 0 && threadIdx.x < sliceheight) y[row] = shared[0][threadIdx.x]; in matmult_seqsell_tiled_kernel9()
137 __global__ void matmultadd_seqsell_tiled_kernel9(PetscInt nrows, PetscInt sliceheight, const PetscI… in matmultadd_seqsell_tiled_kernel9() argument
147 row = slice_id * sliceheight + threadIdx.x % sliceheight; in matmultadd_seqsell_tiled_kernel9()
152 …for (int offset = 16; offset >= sliceheight; offset /= 2) t += __shfl_down_sync(0xffffffff, t, off… in matmultadd_seqsell_tiled_kernel9()
154 if (threadIdx.x < sliceheight) shared[threadIdx.x][threadIdx.y] = t; in matmultadd_seqsell_tiled_kernel9()
156 if (tidy < sliceheight) t = shared[tidy][tidx]; in matmultadd_seqsell_tiled_kernel9()
159 if (tidx == 0 && tidy < sliceheight) shared[0][tidy] = t; in matmultadd_seqsell_tiled_kernel9()
161 …if (row < nrows && threadIdx.y == 0 && threadIdx.x < sliceheight) z[row] = y[row] + shared[0][thre… in matmultadd_seqsell_tiled_kernel9()
185 __global__ void matmult_seqsell_tiled_kernel8(PetscInt nrows, PetscInt sliceheight, PetscInt chunks… in matmult_seqsell_tiled_kernel8() argument
204 … slice_id = start_slice, totalslices = PetscCeilIntMacro(nrows, sliceheight), totalentries … in matmult_seqsell_tiled_kernel8()
208 row = slice_id * sliceheight + threadIdx.x % sliceheight; in matmult_seqsell_tiled_kernel8()
215 row = start_slice * sliceheight + threadIdx.x % sliceheight; in matmult_seqsell_tiled_kernel8()
221 …for (int offset = 16; offset >= sliceheight; offset /= 2) t += __shfl_down_sync(0xffffffff, t, off… in matmult_seqsell_tiled_kernel8()
223 …if (threadIdx.x < sliceheight) shared[threadIdx.x * BLOCKY + threadIdx.y] = t; /* shared[threadIdx… in matmult_seqsell_tiled_kernel8()
225 if (tidy < sliceheight) t = shared[tidy * BLOCKY + tidx]; /* shared[tidy][tidx] */ in matmult_seqsell_tiled_kernel8()
228 if (tidx == 0 && tidy < sliceheight) shared[tidy] = t; /* shared[0][tidy] = t */ in matmult_seqsell_tiled_kernel8()
230 …if (row < nrows && threadIdx.y == 0 && threadIdx.x < sliceheight) atomAdd(y[row], shared[threadIdx… in matmult_seqsell_tiled_kernel8()
240 __global__ void matmultadd_seqsell_tiled_kernel8(PetscInt nrows, PetscInt sliceheight, PetscInt chu… in matmultadd_seqsell_tiled_kernel8() argument
259 … slice_id = start_slice, totalslices = PetscCeilIntMacro(nrows, sliceheight), totalentries … in matmultadd_seqsell_tiled_kernel8()
263 row = slice_id * sliceheight + threadIdx.x % sliceheight; in matmultadd_seqsell_tiled_kernel8()
270 row = start_slice * sliceheight + threadIdx.x % sliceheight; in matmultadd_seqsell_tiled_kernel8()
276 …for (int offset = 16; offset >= sliceheight; offset /= 2) t += __shfl_down_sync(0xffffffff, t, off… in matmultadd_seqsell_tiled_kernel8()
278 …if (threadIdx.x < sliceheight) shared[threadIdx.x * BLOCKY + threadIdx.y] = t; /* shared[threadIdx… in matmultadd_seqsell_tiled_kernel8()
280 if (tidy < sliceheight) t = shared[tidy * BLOCKY + tidx]; /* shared[tidy][tidx] */ in matmultadd_seqsell_tiled_kernel8()
283 if (tidx == 0 && tidy < sliceheight) shared[tidy] = t; /* shared[0][tidy] = t */ in matmultadd_seqsell_tiled_kernel8()
285 …if (row < nrows && threadIdx.y == 0 && threadIdx.x < sliceheight) atomAdd(z[row], shared[threadIdx… in matmultadd_seqsell_tiled_kernel8()
294 static __global__ void matmult_seqsell_tiled_kernel7(PetscInt nrows, PetscInt sliceheight, const Pe… in matmult_seqsell_tiled_kernel7() argument
298 row = slice_id * sliceheight + threadIdx.x % sliceheight; in matmult_seqsell_tiled_kernel7()
304 …for (int offset = 16; offset >= sliceheight; offset /= 2) t += __shfl_down_sync(0xffffffff, t, off… in matmult_seqsell_tiled_kernel7()
305 if (row < nrows && threadIdx.x < sliceheight) y[row] = t; in matmult_seqsell_tiled_kernel7()
309 static __global__ void matmultadd_seqsell_tiled_kernel7(PetscInt nrows, PetscInt sliceheight, const… in matmultadd_seqsell_tiled_kernel7() argument
313 row = slice_id * sliceheight + threadIdx.x % sliceheight; in matmultadd_seqsell_tiled_kernel7()
319 …for (int offset = 16; offset >= sliceheight; offset /= 2) t += __shfl_down_sync(0xffffffff, t, off… in matmultadd_seqsell_tiled_kernel7()
320 if (row < nrows && threadIdx.x < sliceheight) z[row] = y[row] + t; in matmultadd_seqsell_tiled_kernel7()
572 PetscInt nrows = A->rmap->n, sliceheight = a->sliceheight; in MatMult_SeqSELLCUDA() local
584 …sliceheight == 0, PETSC_COMM_SELF, PETSC_ERR_SUP, "The kernel requires a slice height be a divisor… in MatMult_SeqSELLCUDA()
585 …sliceheight != SLICE_HEIGHT), PETSC_COMM_SELF, PETSC_ERR_ARG_OUTOFRANGE, "Kernel choices {2-6} req… in MatMult_SeqSELLCUDA()
599 nblocks = 1 + (nrows - 1) / sliceheight; in MatMult_SeqSELLCUDA()
601 …matmult_seqsell_tiled_kernel9<2><<<nblocks, dim3(32, 2)>>>(nrows, sliceheight, acolidx, aval, slii… in MatMult_SeqSELLCUDA()
603 …matmult_seqsell_tiled_kernel9<4><<<nblocks, dim3(32, 4)>>>(nrows, sliceheight, acolidx, aval, slii… in MatMult_SeqSELLCUDA()
605 …matmult_seqsell_tiled_kernel9<8><<<nblocks, dim3(32, 8)>>>(nrows, sliceheight, acolidx, aval, slii… in MatMult_SeqSELLCUDA()
607 …matmult_seqsell_tiled_kernel9<16><<<nblocks, dim3(32, 16)>>>(nrows, sliceheight, acolidx, aval, sl… in MatMult_SeqSELLCUDA()
609 …matmult_seqsell_tiled_kernel9<32><<<nblocks, dim3(32, 32)>>>(nrows, sliceheight, acolidx, aval, sl… in MatMult_SeqSELLCUDA()
611 …matmult_seqsell_tiled_kernel9<2><<<nblocks, dim3(32, 2)>>>(nrows, sliceheight, acolidx, aval, slii… in MatMult_SeqSELLCUDA()
615 nblocks = 1 + (nrows - 1) / (2 * sliceheight); in MatMult_SeqSELLCUDA()
617 …matmult_seqsell_tiled_kernel7<<<nblocks, dim3(32, 2)>>>(nrows, sliceheight, acolidx, aval, sliidx,… in MatMult_SeqSELLCUDA()
619 …matmult_seqsell_tiled_kernel7<<<nblocks, dim3(32, 4)>>>(nrows, sliceheight, acolidx, aval, sliidx,… in MatMult_SeqSELLCUDA()
621 …matmult_seqsell_tiled_kernel7<<<nblocks, dim3(32, 8)>>>(nrows, sliceheight, acolidx, aval, sliidx,… in MatMult_SeqSELLCUDA()
623 …matmult_seqsell_tiled_kernel7<<<nblocks, dim3(32, 16)>>>(nrows, sliceheight, acolidx, aval, sliidx… in MatMult_SeqSELLCUDA()
625 …matmult_seqsell_tiled_kernel7<<<nblocks, dim3(32, 32)>>>(nrows, sliceheight, acolidx, aval, sliidx… in MatMult_SeqSELLCUDA()
627 …matmult_seqsell_tiled_kernel7<<<nblocks, dim3(32, 2)>>>(nrows, sliceheight, acolidx, aval, sliidx,… in MatMult_SeqSELLCUDA()
653 …matmult_seqsell_basic_kernel<<<nblocks, blocksize>>>(nrows, sliceheight, acolidx, aval, sliidx, x,… in MatMult_SeqSELLCUDA()
666 …matmult_seqsell_tiled_kernel8<2><<<nblocks, dim3(32, 2)>>>(nrows, sliceheight, chunksperblock, nch… in MatMult_SeqSELLCUDA()
668 …matmult_seqsell_tiled_kernel8<4><<<nblocks, dim3(32, 4)>>>(nrows, sliceheight, chunksperblock, nch… in MatMult_SeqSELLCUDA()
670 …matmult_seqsell_tiled_kernel8<8><<<nblocks, dim3(32, 8)>>>(nrows, sliceheight, chunksperblock, nch… in MatMult_SeqSELLCUDA()
672 …matmult_seqsell_tiled_kernel8<16><<<nblocks, dim3(32, 16)>>>(nrows, sliceheight, chunksperblock, n… in MatMult_SeqSELLCUDA()
674 …matmult_seqsell_tiled_kernel8<32><<<nblocks, dim3(32, 32)>>>(nrows, sliceheight, chunksperblock, n… in MatMult_SeqSELLCUDA()
676 …matmult_seqsell_tiled_kernel8<2><<<nblocks, dim3(32, 2)>>>(nrows, sliceheight, chunksperblock, nch… in MatMult_SeqSELLCUDA()
679 PetscInt avgslicesize = sliceheight * a->avgslicewidth; in MatMult_SeqSELLCUDA()
681 if (sliceheight * a->maxslicewidth < 2048 && nrows > 100000) { in MatMult_SeqSELLCUDA()
682 nblocks = 1 + (nrows - 1) / (2 * sliceheight); /* two slices per block */ in MatMult_SeqSELLCUDA()
683 …matmult_seqsell_tiled_kernel7<<<nblocks, dim3(32, 2)>>>(nrows, sliceheight, acolidx, aval, sliidx,… in MatMult_SeqSELLCUDA()
685 nblocks = 1 + (nrows - 1) / sliceheight; in MatMult_SeqSELLCUDA()
686 …matmult_seqsell_tiled_kernel9<2><<<nblocks, dim3(32, 2)>>>(nrows, sliceheight, acolidx, aval, slii… in MatMult_SeqSELLCUDA()
689 nblocks = 1 + (nrows - 1) / sliceheight; in MatMult_SeqSELLCUDA()
690 …matmult_seqsell_tiled_kernel9<8><<<nblocks, dim3(32, 8)>>>(nrows, sliceheight, acolidx, aval, slii… in MatMult_SeqSELLCUDA()
692 nblocks = 1 + (nrows - 1) / sliceheight; in MatMult_SeqSELLCUDA()
693 …matmult_seqsell_tiled_kernel9<16><<<nblocks, dim3(32, 16)>>>(nrows, sliceheight, acolidx, aval, sl… in MatMult_SeqSELLCUDA()
714 PetscInt nrows = A->rmap->n, sliceheight = a->sliceheight; in MatMultAdd_SeqSELLCUDA() local
725 …sliceheight == 0, PETSC_COMM_SELF, PETSC_ERR_SUP, "The kernel requires a slice height be a divisor… in MatMultAdd_SeqSELLCUDA()
726 …sliceheight != SLICE_HEIGHT), PETSC_COMM_SELF, PETSC_ERR_ARG_OUTOFRANGE, "Kernel choices {2-6} req… in MatMultAdd_SeqSELLCUDA()
739 nblocks = 1 + (nrows - 1) / sliceheight; in MatMultAdd_SeqSELLCUDA()
741 …matmultadd_seqsell_tiled_kernel9<2><<<nblocks, dim3(32, 2)>>>(nrows, sliceheight, acolidx, aval, s… in MatMultAdd_SeqSELLCUDA()
743 …matmultadd_seqsell_tiled_kernel9<4><<<nblocks, dim3(32, 4)>>>(nrows, sliceheight, acolidx, aval, s… in MatMultAdd_SeqSELLCUDA()
745 …matmultadd_seqsell_tiled_kernel9<8><<<nblocks, dim3(32, 8)>>>(nrows, sliceheight, acolidx, aval, s… in MatMultAdd_SeqSELLCUDA()
747 …matmultadd_seqsell_tiled_kernel9<16><<<nblocks, dim3(32, 16)>>>(nrows, sliceheight, acolidx, aval,… in MatMultAdd_SeqSELLCUDA()
749 …matmultadd_seqsell_tiled_kernel9<32><<<nblocks, dim3(32, 32)>>>(nrows, sliceheight, acolidx, aval,… in MatMultAdd_SeqSELLCUDA()
751 …matmultadd_seqsell_tiled_kernel9<2><<<nblocks, dim3(32, 2)>>>(nrows, sliceheight, acolidx, aval, s… in MatMultAdd_SeqSELLCUDA()
762 …matmultadd_seqsell_tiled_kernel8<2><<<nblocks, dim3(32, 2)>>>(nrows, sliceheight, chunksperblock, … in MatMultAdd_SeqSELLCUDA()
764 …matmultadd_seqsell_tiled_kernel8<4><<<nblocks, dim3(32, 4)>>>(nrows, sliceheight, chunksperblock, … in MatMultAdd_SeqSELLCUDA()
766 …matmultadd_seqsell_tiled_kernel8<8><<<nblocks, dim3(32, 8)>>>(nrows, sliceheight, chunksperblock, … in MatMultAdd_SeqSELLCUDA()
768 …matmultadd_seqsell_tiled_kernel8<16><<<nblocks, dim3(32, 16)>>>(nrows, sliceheight, chunksperblock… in MatMultAdd_SeqSELLCUDA()
770 …matmultadd_seqsell_tiled_kernel8<32><<<nblocks, dim3(32, 32)>>>(nrows, sliceheight, chunksperblock… in MatMultAdd_SeqSELLCUDA()
772 …matmultadd_seqsell_tiled_kernel8<2><<<nblocks, dim3(32, 2)>>>(nrows, sliceheight, chunksperblock, … in MatMultAdd_SeqSELLCUDA()
776 nblocks = 1 + (nrows - 1) / (2 * sliceheight); in MatMultAdd_SeqSELLCUDA()
778 …matmultadd_seqsell_tiled_kernel7<<<nblocks, dim3(32, 2)>>>(nrows, sliceheight, acolidx, aval, slii… in MatMultAdd_SeqSELLCUDA()
780 …matmultadd_seqsell_tiled_kernel7<<<nblocks, dim3(32, 4)>>>(nrows, sliceheight, acolidx, aval, slii… in MatMultAdd_SeqSELLCUDA()
782 …matmultadd_seqsell_tiled_kernel7<<<nblocks, dim3(32, 8)>>>(nrows, sliceheight, acolidx, aval, slii… in MatMultAdd_SeqSELLCUDA()
784 …matmultadd_seqsell_tiled_kernel7<<<nblocks, dim3(32, 16)>>>(nrows, sliceheight, acolidx, aval, sli… in MatMultAdd_SeqSELLCUDA()
786 …matmultadd_seqsell_tiled_kernel7<<<nblocks, dim3(32, 32)>>>(nrows, sliceheight, acolidx, aval, sli… in MatMultAdd_SeqSELLCUDA()
788 …matmultadd_seqsell_tiled_kernel7<<<nblocks, dim3(32, 2)>>>(nrows, sliceheight, acolidx, aval, slii… in MatMultAdd_SeqSELLCUDA()
814 …matmultadd_seqsell_basic_kernel<<<nblocks, blocksize>>>(nrows, sliceheight, acolidx, aval, sliidx,… in MatMultAdd_SeqSELLCUDA()
827 …matmultadd_seqsell_tiled_kernel8<2><<<nblocks, dim3(32, 2)>>>(nrows, sliceheight, chunksperblock, … in MatMultAdd_SeqSELLCUDA()
829 …matmultadd_seqsell_tiled_kernel8<4><<<nblocks, dim3(32, 4)>>>(nrows, sliceheight, chunksperblock, … in MatMultAdd_SeqSELLCUDA()
831 …matmultadd_seqsell_tiled_kernel8<8><<<nblocks, dim3(32, 8)>>>(nrows, sliceheight, chunksperblock, … in MatMultAdd_SeqSELLCUDA()
833 …matmultadd_seqsell_tiled_kernel8<16><<<nblocks, dim3(32, 16)>>>(nrows, sliceheight, chunksperblock… in MatMultAdd_SeqSELLCUDA()
835 …matmultadd_seqsell_tiled_kernel8<32><<<nblocks, dim3(32, 32)>>>(nrows, sliceheight, chunksperblock… in MatMultAdd_SeqSELLCUDA()
837 …matmultadd_seqsell_tiled_kernel8<2><<<nblocks, dim3(32, 2)>>>(nrows, sliceheight, chunksperblock, … in MatMultAdd_SeqSELLCUDA()
840 PetscInt avgslicesize = sliceheight * a->avgslicewidth; in MatMultAdd_SeqSELLCUDA()
842 if (sliceheight * a->maxslicewidth < 2048 && nrows > 100000) { in MatMultAdd_SeqSELLCUDA()
843 nblocks = 1 + (nrows - 1) / (2 * sliceheight); /* two slices per block */ in MatMultAdd_SeqSELLCUDA()
844 …matmultadd_seqsell_tiled_kernel7<<<nblocks, dim3(32, 2)>>>(nrows, sliceheight, acolidx, aval, slii… in MatMultAdd_SeqSELLCUDA()
846 nblocks = 1 + (nrows - 1) / sliceheight; in MatMultAdd_SeqSELLCUDA()
847 …matmultadd_seqsell_tiled_kernel9<2><<<nblocks, dim3(32, 2)>>>(nrows, sliceheight, acolidx, aval, s… in MatMultAdd_SeqSELLCUDA()
850 nblocks = 1 + (nrows - 1) / sliceheight; in MatMultAdd_SeqSELLCUDA()
851 …matmultadd_seqsell_tiled_kernel9<8><<<nblocks, dim3(32, 8)>>>(nrows, sliceheight, acolidx, aval, s… in MatMultAdd_SeqSELLCUDA()
853 nblocks = 1 + (nrows - 1) / sliceheight; in MatMultAdd_SeqSELLCUDA()
854 …matmultadd_seqsell_tiled_kernel9<16><<<nblocks, dim3(32, 16)>>>(nrows, sliceheight, acolidx, aval,… in MatMultAdd_SeqSELLCUDA()