seq/seqhip/sellhip.hip.cxx

581   PetscInt           nblocks, blocksize = 512; /* blocksize is fixed to be 512 */  in MatMult_SeqSELLHIP()  local
604     nblocks = 1 + (nrows - 1) / sliceheight;  in MatMult_SeqSELLHIP()
606 …matmult_seqsell_tiled_kernel9<2><<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, acolidx, ava…  in MatMult_SeqSELLHIP()
608 …matmult_seqsell_tiled_kernel9<4><<<nblocks, dim3(WARP_SIZE, 4)>>>(nrows, sliceheight, acolidx, ava…  in MatMult_SeqSELLHIP()
610 …matmult_seqsell_tiled_kernel9<8><<<nblocks, dim3(WARP_SIZE, 8)>>>(nrows, sliceheight, acolidx, ava…  in MatMult_SeqSELLHIP()
612 …matmult_seqsell_tiled_kernel9<16><<<nblocks, dim3(WARP_SIZE, 16)>>>(nrows, sliceheight, acolidx, a…  in MatMult_SeqSELLHIP()
614 …matmult_seqsell_tiled_kernel9<2><<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, acolidx, ava…  in MatMult_SeqSELLHIP()
618     nblocks = 1 + (nrows - 1) / (hipstruct->blocky * sliceheight);  in MatMult_SeqSELLHIP()
620 …matmult_seqsell_tiled_kernel7<<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, acolidx, aval, …  in MatMult_SeqSELLHIP()
622 …matmult_seqsell_tiled_kernel7<<<nblocks, dim3(WARP_SIZE, 4)>>>(nrows, sliceheight, acolidx, aval, …  in MatMult_SeqSELLHIP()
624 …matmult_seqsell_tiled_kernel7<<<nblocks, dim3(WARP_SIZE, 8)>>>(nrows, sliceheight, acolidx, aval, …  in MatMult_SeqSELLHIP()
626 …matmult_seqsell_tiled_kernel7<<<nblocks, dim3(WARP_SIZE, 16)>>>(nrows, sliceheight, acolidx, aval,…  in MatMult_SeqSELLHIP()
628       nblocks = 1 + (nrows - 1) / (2 * sliceheight);  in MatMult_SeqSELLHIP()
629 …matmult_seqsell_tiled_kernel7<<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, acolidx, aval, …  in MatMult_SeqSELLHIP()
634     nblocks = 1 + (nrows - 1) / (blocksize / 32); /* 1 slice per block if sliceheight=32 */  in MatMult_SeqSELLHIP()
635 …matmult_seqsell_tiled_kernel6<<<nblocks, block32>>>(nrows, sliceheight, acolidx, aval, sliidx, x, …  in MatMult_SeqSELLHIP()
638     nblocks = 1 + (nrows - 1) / (blocksize / 16); /* 2 slices per block if sliceheight=32*/  in MatMult_SeqSELLHIP()
639 …matmult_seqsell_tiled_kernel5<<<nblocks, block16>>>(nrows, sliceheight, acolidx, aval, sliidx, x, …  in MatMult_SeqSELLHIP()
642     nblocks = 1 + (nrows - 1) / (blocksize / 8); /* 4 slices per block if sliceheight=32 */  in MatMult_SeqSELLHIP()
643 …matmult_seqsell_tiled_kernel4<<<nblocks, block8>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y…  in MatMult_SeqSELLHIP()
646     nblocks = 1 + (nrows - 1) / (blocksize / 4); /* 8 slices per block if sliceheight=32 */  in MatMult_SeqSELLHIP()
647 …matmult_seqsell_tiled_kernel3<<<nblocks, block4>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y…  in MatMult_SeqSELLHIP()
650     nblocks = 1 + (nrows - 1) / (blocksize / 2);  in MatMult_SeqSELLHIP()
651 …matmult_seqsell_tiled_kernel2<<<nblocks, block2>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y…  in MatMult_SeqSELLHIP()
654     nblocks = 1 + (nrows - 1) / blocksize;  in MatMult_SeqSELLHIP()
655 …matmult_seqsell_basic_kernel<<<nblocks, blocksize>>>(nrows, sliceheight, acolidx, aval, sliidx, x,…  in MatMult_SeqSELLHIP()
665       nblocks         = 1 + (nchunks - 1) / chunksperblock;  in MatMult_SeqSELLHIP()
668 …matmult_seqsell_tiled_kernel8<2><<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, chunksperblo…  in MatMult_SeqSELLHIP()
670 …matmult_seqsell_tiled_kernel8<4><<<nblocks, dim3(WARP_SIZE, 4)>>>(nrows, sliceheight, chunksperblo…  in MatMult_SeqSELLHIP()
672 …matmult_seqsell_tiled_kernel8<8><<<nblocks, dim3(WARP_SIZE, 8)>>>(nrows, sliceheight, chunksperblo…  in MatMult_SeqSELLHIP()
674 …matmult_seqsell_tiled_kernel8<16><<<nblocks, dim3(WARP_SIZE, 16)>>>(nrows, sliceheight, chunksperb…  in MatMult_SeqSELLHIP()
676 …matmult_seqsell_tiled_kernel8<2><<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, chunksperblo…  in MatMult_SeqSELLHIP()
682           nblocks = 1 + (nrows - 1) / (2 * sliceheight); /* two slices per block */  in MatMult_SeqSELLHIP()
683 …matmult_seqsell_tiled_kernel7<<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, acolidx, aval, …  in MatMult_SeqSELLHIP()
685           nblocks = 1 + (nrows - 1) / sliceheight;  in MatMult_SeqSELLHIP()
686 …matmult_seqsell_tiled_kernel9<2><<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, acolidx, ava…  in MatMult_SeqSELLHIP()
689         nblocks = 1 + (nrows - 1) / sliceheight;  in MatMult_SeqSELLHIP()
690 …matmult_seqsell_tiled_kernel9<8><<<nblocks, dim3(WARP_SIZE, 8)>>>(nrows, sliceheight, acolidx, ava…  in MatMult_SeqSELLHIP()
692         nblocks = 1 + (nrows - 1) / sliceheight;  in MatMult_SeqSELLHIP()
693 …matmult_seqsell_tiled_kernel9<16><<<nblocks, dim3(WARP_SIZE, 16)>>>(nrows, sliceheight, acolidx, a…  in MatMult_SeqSELLHIP()
729     PetscInt nblocks, blocksize = 512;  in MatMultAdd_SeqSELLHIP()  local
739       nblocks = 1 + (nrows - 1) / sliceheight;  in MatMultAdd_SeqSELLHIP()
741 …matmultadd_seqsell_tiled_kernel9<2><<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, acolidx, …  in MatMultAdd_SeqSELLHIP()
743 …matmultadd_seqsell_tiled_kernel9<4><<<nblocks, dim3(WARP_SIZE, 4)>>>(nrows, sliceheight, acolidx, …  in MatMultAdd_SeqSELLHIP()
745 …matmultadd_seqsell_tiled_kernel9<8><<<nblocks, dim3(WARP_SIZE, 8)>>>(nrows, sliceheight, acolidx, …  in MatMultAdd_SeqSELLHIP()
747 …matmultadd_seqsell_tiled_kernel9<16><<<nblocks, dim3(WARP_SIZE, 16)>>>(nrows, sliceheight, acolidx…  in MatMultAdd_SeqSELLHIP()
749 …matmultadd_seqsell_tiled_kernel9<2><<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, acolidx, …  in MatMultAdd_SeqSELLHIP()
757       nblocks         = 1 + (nchunks - 1) / chunksperblock;  in MatMultAdd_SeqSELLHIP()
760 …matmultadd_seqsell_tiled_kernel8<2><<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, chunksper…  in MatMultAdd_SeqSELLHIP()
762 …matmultadd_seqsell_tiled_kernel8<4><<<nblocks, dim3(WARP_SIZE, 4)>>>(nrows, sliceheight, chunksper…  in MatMultAdd_SeqSELLHIP()
764 …matmultadd_seqsell_tiled_kernel8<8><<<nblocks, dim3(WARP_SIZE, 8)>>>(nrows, sliceheight, chunksper…  in MatMultAdd_SeqSELLHIP()
766 …matmultadd_seqsell_tiled_kernel8<16><<<nblocks, dim3(WARP_SIZE, 16)>>>(nrows, sliceheight, chunksp…  in MatMultAdd_SeqSELLHIP()
768 …matmultadd_seqsell_tiled_kernel8<2><<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, chunksper…  in MatMultAdd_SeqSELLHIP()
772       nblocks = 1 + (nrows - 1) / (blocky * sliceheight);  in MatMultAdd_SeqSELLHIP()
774 …matmultadd_seqsell_tiled_kernel7<<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, acolidx, ava…  in MatMultAdd_SeqSELLHIP()
776 …matmultadd_seqsell_tiled_kernel7<<<nblocks, dim3(WARP_SIZE, 4)>>>(nrows, sliceheight, acolidx, ava…  in MatMultAdd_SeqSELLHIP()
778 …matmultadd_seqsell_tiled_kernel7<<<nblocks, dim3(WARP_SIZE, 8)>>>(nrows, sliceheight, acolidx, ava…  in MatMultAdd_SeqSELLHIP()
780 …matmultadd_seqsell_tiled_kernel7<<<nblocks, dim3(WARP_SIZE, 16)>>>(nrows, sliceheight, acolidx, av…  in MatMultAdd_SeqSELLHIP()
782         nblocks = 1 + (nrows - 1) / (2 * sliceheight);  in MatMultAdd_SeqSELLHIP()
783 …matmultadd_seqsell_tiled_kernel7<<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, acolidx, ava…  in MatMultAdd_SeqSELLHIP()
788       nblocks = 1 + (nrows - 1) / (blocksize / 32);  in MatMultAdd_SeqSELLHIP()
789 …matmultadd_seqsell_tiled_kernel6<<<nblocks, block32>>>(nrows, sliceheight, acolidx, aval, sliidx, …  in MatMultAdd_SeqSELLHIP()
792       nblocks = 1 + (nrows - 1) / (blocksize / 16);  in MatMultAdd_SeqSELLHIP()
793 …matmultadd_seqsell_tiled_kernel5<<<nblocks, block16>>>(nrows, sliceheight, acolidx, aval, sliidx, …  in MatMultAdd_SeqSELLHIP()
796       nblocks = 1 + (nrows - 1) / (blocksize / 8);  in MatMultAdd_SeqSELLHIP()
797 …matmultadd_seqsell_tiled_kernel4<<<nblocks, block8>>>(nrows, sliceheight, acolidx, aval, sliidx, x…  in MatMultAdd_SeqSELLHIP()
800       nblocks = 1 + (nrows - 1) / (blocksize / 4);  in MatMultAdd_SeqSELLHIP()
801 …matmultadd_seqsell_tiled_kernel3<<<nblocks, block4>>>(nrows, sliceheight, acolidx, aval, sliidx, x…  in MatMultAdd_SeqSELLHIP()
804       nblocks = 1 + (nrows - 1) / (blocksize / 2);  in MatMultAdd_SeqSELLHIP()
805 …matmultadd_seqsell_tiled_kernel2<<<nblocks, block2>>>(nrows, sliceheight, acolidx, aval, sliidx, x…  in MatMultAdd_SeqSELLHIP()
808       nblocks = 1 + (nrows - 1) / blocksize;  in MatMultAdd_SeqSELLHIP()
809 …matmultadd_seqsell_basic_kernel<<<nblocks, blocksize>>>(nrows, sliceheight, acolidx, aval, sliidx,…  in MatMultAdd_SeqSELLHIP()
819         nblocks         = 1 + (nchunks - 1) / chunksperblock;  in MatMultAdd_SeqSELLHIP()
822 …matmultadd_seqsell_tiled_kernel8<2><<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, chunksper…  in MatMultAdd_SeqSELLHIP()
824 …matmultadd_seqsell_tiled_kernel8<4><<<nblocks, dim3(WARP_SIZE, 4)>>>(nrows, sliceheight, chunksper…  in MatMultAdd_SeqSELLHIP()
826 …matmultadd_seqsell_tiled_kernel8<8><<<nblocks, dim3(WARP_SIZE, 8)>>>(nrows, sliceheight, chunksper…  in MatMultAdd_SeqSELLHIP()
828 …matmultadd_seqsell_tiled_kernel8<16><<<nblocks, dim3(WARP_SIZE, 16)>>>(nrows, sliceheight, chunksp…  in MatMultAdd_SeqSELLHIP()
830 …matmultadd_seqsell_tiled_kernel8<2><<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, chunksper…  in MatMultAdd_SeqSELLHIP()
836             nblocks = 1 + (nrows - 1) / (2 * sliceheight); /* two slices per block */  in MatMultAdd_SeqSELLHIP()
837 …matmultadd_seqsell_tiled_kernel7<<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, acolidx, ava…  in MatMultAdd_SeqSELLHIP()
839             nblocks = 1 + (nrows - 1) / sliceheight;  in MatMultAdd_SeqSELLHIP()
840 …matmultadd_seqsell_tiled_kernel9<2><<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, acolidx, …  in MatMultAdd_SeqSELLHIP()
843           nblocks = 1 + (nrows - 1) / sliceheight;  in MatMultAdd_SeqSELLHIP()
844 …matmultadd_seqsell_tiled_kernel9<8><<<nblocks, dim3(WARP_SIZE, 8)>>>(nrows, sliceheight, acolidx, …  in MatMultAdd_SeqSELLHIP()
846           nblocks = 1 + (nrows - 1) / sliceheight;  in MatMultAdd_SeqSELLHIP()
847 …matmultadd_seqsell_tiled_kernel9<16><<<nblocks, dim3(WARP_SIZE, 16)>>>(nrows, sliceheight, acolidx…  in MatMultAdd_SeqSELLHIP()