Lines Matching refs:acolidx
75 …qsell_basic_kernel(PetscInt nrows, PetscInt sliceheight, const PetscInt *acolidx, const MatScalar … in matmult_seqsell_basic_kernel() argument
85 …ice_id] + row_in_slice; i < sliidx[slice_id + 1]; i += sliceheight) sum += aval[i] * x[acolidx[i]]; in matmult_seqsell_basic_kernel()
90 …qsell_basic_kernel(PetscInt nrows, PetscInt sliceheight, const PetscInt *acolidx, const MatScalar … in matmultadd_seqsell_basic_kernel() argument
100 …ice_id] + row_in_slice; i < sliidx[slice_id + 1]; i += sliceheight) sum += aval[i] * x[acolidx[i]]; in matmultadd_seqsell_basic_kernel()
108 …sell_tiled_kernel9(PetscInt nrows, PetscInt sliceheight, const PetscInt *acolidx, const MatScalar … in matmult_seqsell_tiled_kernel9() argument
120 …Idx.x + 32 * threadIdx.y; i < sliidx[slice_id + 1]; i += 32 * BLOCKY) t += aval[i] * x[acolidx[i]]; in matmult_seqsell_tiled_kernel9()
137 …sell_tiled_kernel9(PetscInt nrows, PetscInt sliceheight, const PetscInt *acolidx, const MatScalar … in matmultadd_seqsell_tiled_kernel9() argument
149 …Idx.x + 32 * threadIdx.y; i < sliidx[slice_id + 1]; i += 32 * BLOCKY) t += aval[i] * x[acolidx[i]]; in matmultadd_seqsell_tiled_kernel9()
185 …, PetscInt totalchunks, const PetscInt *chunk_slice_map, const PetscInt *acolidx, const MatScalar … in matmult_seqsell_tiled_kernel8() argument
209 if (row < nrows && gid < totalentries) t = aval[gid] * x[acolidx[gid]]; in matmult_seqsell_tiled_kernel8()
216 if (row < nrows) t += aval[gid] * x[acolidx[gid]]; in matmult_seqsell_tiled_kernel8()
240 …, PetscInt totalchunks, const PetscInt *chunk_slice_map, const PetscInt *acolidx, const MatScalar … in matmultadd_seqsell_tiled_kernel8() argument
264 if (row < nrows && gid < totalentries) t = aval[gid] * x[acolidx[gid]]; in matmultadd_seqsell_tiled_kernel8()
271 if (row < nrows) t += aval[gid] * x[acolidx[gid]]; in matmultadd_seqsell_tiled_kernel8()
294 …sell_tiled_kernel7(PetscInt nrows, PetscInt sliceheight, const PetscInt *acolidx, const MatScalar … in matmult_seqsell_tiled_kernel7() argument
301 … = sliidx[slice_id] + threadIdx.x; i < sliidx[slice_id + 1]; i += 32) t += aval[i] * x[acolidx[i]]; in matmult_seqsell_tiled_kernel7()
309 …sell_tiled_kernel7(PetscInt nrows, PetscInt sliceheight, const PetscInt *acolidx, const MatScalar … in matmultadd_seqsell_tiled_kernel7() argument
316 … = sliidx[slice_id] + threadIdx.x; i < sliidx[slice_id + 1]; i += 32) t += aval[i] * x[acolidx[i]]; in matmultadd_seqsell_tiled_kernel7()
326 static __global__ void matmult_seqsell_tiled_kernel6(PetscInt nrows, const PetscInt *acolidx, const… in matmult_seqsell_tiled_kernel6() argument
337 …ICE_HEIGHT * blockDim.y) shared[threadIdx.y * blockDim.x + threadIdx.x] += aval[i] * x[acolidx[i]]; in matmult_seqsell_tiled_kernel6()
354 static __global__ void matmult_seqsell_tiled_kernel5(PetscInt nrows, const PetscInt *acolidx, const… in matmult_seqsell_tiled_kernel5() argument
365 …ICE_HEIGHT * blockDim.y) shared[threadIdx.y * blockDim.x + threadIdx.x] += aval[i] * x[acolidx[i]]; in matmult_seqsell_tiled_kernel5()
380 static __global__ void matmult_seqsell_tiled_kernel4(PetscInt nrows, const PetscInt *acolidx, const… in matmult_seqsell_tiled_kernel4() argument
391 …ICE_HEIGHT * blockDim.y) shared[threadIdx.y * blockDim.x + threadIdx.x] += aval[i] * x[acolidx[i]]; in matmult_seqsell_tiled_kernel4()
404 static __global__ void matmult_seqsell_tiled_kernel3(PetscInt nrows, const PetscInt *acolidx, const… in matmult_seqsell_tiled_kernel3() argument
415 …ICE_HEIGHT * blockDim.y) shared[threadIdx.y * blockDim.x + threadIdx.x] += aval[i] * x[acolidx[i]]; in matmult_seqsell_tiled_kernel3()
426 static __global__ void matmult_seqsell_tiled_kernel2(PetscInt nrows, const PetscInt *acolidx, const… in matmult_seqsell_tiled_kernel2() argument
437 …ICE_HEIGHT * blockDim.y) shared[threadIdx.y * blockDim.x + threadIdx.x] += aval[i] * x[acolidx[i]]; in matmult_seqsell_tiled_kernel2()
446 static __global__ void matmultadd_seqsell_tiled_kernel6(PetscInt nrows, const PetscInt *acolidx, co… in matmultadd_seqsell_tiled_kernel6() argument
457 …ICE_HEIGHT * blockDim.y) shared[threadIdx.y * blockDim.x + threadIdx.x] += aval[i] * x[acolidx[i]]; in matmultadd_seqsell_tiled_kernel6()
474 static __global__ void matmultadd_seqsell_tiled_kernel5(PetscInt nrows, const PetscInt *acolidx, co… in matmultadd_seqsell_tiled_kernel5() argument
485 …ICE_HEIGHT * blockDim.y) shared[threadIdx.y * blockDim.x + threadIdx.x] += aval[i] * x[acolidx[i]]; in matmultadd_seqsell_tiled_kernel5()
500 static __global__ void matmultadd_seqsell_tiled_kernel4(PetscInt nrows, const PetscInt *acolidx, co… in matmultadd_seqsell_tiled_kernel4() argument
511 …ICE_HEIGHT * blockDim.y) shared[threadIdx.y * blockDim.x + threadIdx.x] += aval[i] * x[acolidx[i]]; in matmultadd_seqsell_tiled_kernel4()
524 static __global__ void matmultadd_seqsell_tiled_kernel3(PetscInt nrows, const PetscInt *acolidx, co… in matmultadd_seqsell_tiled_kernel3() argument
535 …ICE_HEIGHT * blockDim.y) shared[threadIdx.y * blockDim.x + threadIdx.x] += aval[i] * x[acolidx[i]]; in matmultadd_seqsell_tiled_kernel3()
546 static __global__ void matmultadd_seqsell_tiled_kernel2(PetscInt nrows, const PetscInt *acolidx, co… in matmultadd_seqsell_tiled_kernel2() argument
557 …ICE_HEIGHT * blockDim.y) shared[threadIdx.y * blockDim.x + threadIdx.x] += aval[i] * x[acolidx[i]]; in matmultadd_seqsell_tiled_kernel2()
574 PetscInt *acolidx; in MatMult_SeqSELLCUDA() local
589 acolidx = cudastruct->colidx; 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()
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()
633 matmult_seqsell_tiled_kernel6<<<nblocks, block32>>>(nrows, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLCUDA()
637 matmult_seqsell_tiled_kernel5<<<nblocks, block16>>>(nrows, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLCUDA()
641 matmult_seqsell_tiled_kernel4<<<nblocks, block8>>>(nrows, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLCUDA()
645 matmult_seqsell_tiled_kernel3<<<nblocks, block4>>>(nrows, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLCUDA()
649 matmult_seqsell_tiled_kernel2<<<nblocks, block2>>>(nrows, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLCUDA()
653 …matmult_seqsell_basic_kernel<<<nblocks, blocksize>>>(nrows, sliceheight, acolidx, aval, sliidx, x,… in MatMult_SeqSELLCUDA()
666 …(32, 2)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x,… in MatMult_SeqSELLCUDA()
668 …(32, 4)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x,… in MatMult_SeqSELLCUDA()
670 …(32, 8)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x,… in MatMult_SeqSELLCUDA()
672 …32, 16)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x,… in MatMult_SeqSELLCUDA()
674 …32, 32)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x,… in MatMult_SeqSELLCUDA()
676 …(32, 2)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x,… in MatMult_SeqSELLCUDA()
683 …matmult_seqsell_tiled_kernel7<<<nblocks, dim3(32, 2)>>>(nrows, sliceheight, acolidx, aval, sliidx,… in MatMult_SeqSELLCUDA()
686 …matmult_seqsell_tiled_kernel9<2><<<nblocks, dim3(32, 2)>>>(nrows, sliceheight, acolidx, aval, slii… in MatMult_SeqSELLCUDA()
690 …matmult_seqsell_tiled_kernel9<8><<<nblocks, dim3(32, 8)>>>(nrows, sliceheight, acolidx, aval, slii… in MatMult_SeqSELLCUDA()
693 …matmult_seqsell_tiled_kernel9<16><<<nblocks, dim3(32, 16)>>>(nrows, sliceheight, acolidx, aval, sl… in MatMult_SeqSELLCUDA()
716 PetscInt *acolidx = cudastruct->colidx; in MatMultAdd_SeqSELLCUDA() local
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 …(32, 2)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x,… in MatMultAdd_SeqSELLCUDA()
764 …(32, 4)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x,… in MatMultAdd_SeqSELLCUDA()
766 …(32, 8)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x,… in MatMultAdd_SeqSELLCUDA()
768 …32, 16)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x,… in MatMultAdd_SeqSELLCUDA()
770 …32, 32)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x,… in MatMultAdd_SeqSELLCUDA()
772 …(32, 2)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x,… 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()
794 matmultadd_seqsell_tiled_kernel6<<<nblocks, block32>>>(nrows, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLCUDA()
798 matmultadd_seqsell_tiled_kernel5<<<nblocks, block16>>>(nrows, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLCUDA()
802 matmultadd_seqsell_tiled_kernel4<<<nblocks, block8>>>(nrows, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLCUDA()
806 matmultadd_seqsell_tiled_kernel3<<<nblocks, block4>>>(nrows, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLCUDA()
810 matmultadd_seqsell_tiled_kernel2<<<nblocks, block2>>>(nrows, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLCUDA()
814 …matmultadd_seqsell_basic_kernel<<<nblocks, blocksize>>>(nrows, sliceheight, acolidx, aval, sliidx,… in MatMultAdd_SeqSELLCUDA()
827 …(32, 2)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x,… in MatMultAdd_SeqSELLCUDA()
829 …(32, 4)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x,… in MatMultAdd_SeqSELLCUDA()
831 …(32, 8)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x,… in MatMultAdd_SeqSELLCUDA()
833 …32, 16)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x,… in MatMultAdd_SeqSELLCUDA()
835 …32, 32)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x,… in MatMultAdd_SeqSELLCUDA()
837 …(32, 2)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x,… in MatMultAdd_SeqSELLCUDA()
844 …matmultadd_seqsell_tiled_kernel7<<<nblocks, dim3(32, 2)>>>(nrows, sliceheight, acolidx, aval, slii… in MatMultAdd_SeqSELLCUDA()
847 …d_seqsell_tiled_kernel9<2><<<nblocks, dim3(32, 2)>>>(nrows, sliceheight, acolidx, aval, sliidx, x,… in MatMultAdd_SeqSELLCUDA()
851 …matmultadd_seqsell_tiled_kernel9<8><<<nblocks, dim3(32, 8)>>>(nrows, sliceheight, acolidx, aval, s… in MatMultAdd_SeqSELLCUDA()
854 …seqsell_tiled_kernel9<16><<<nblocks, dim3(32, 16)>>>(nrows, sliceheight, acolidx, aval, sliidx, x,… in MatMultAdd_SeqSELLCUDA()