5 #define SLICE_HEIGHT 16  macro
333     slice_id     = row / SLICE_HEIGHT;  in matmult_seqsell_tiled_kernel6()
334     row_in_slice = row % SLICE_HEIGHT;  in matmult_seqsell_tiled_kernel6()
337 …r (i = sliidx[slice_id] + row_in_slice + SLICE_HEIGHT * threadIdx.y; i < sliidx[slice_id + 1]; i +…  in matmult_seqsell_tiled_kernel6()
361     slice_id     = row / SLICE_HEIGHT;  in matmult_seqsell_tiled_kernel5()
362     row_in_slice = row % SLICE_HEIGHT;  in matmult_seqsell_tiled_kernel5()
365 …r (i = sliidx[slice_id] + row_in_slice + SLICE_HEIGHT * threadIdx.y; i < sliidx[slice_id + 1]; i +…  in matmult_seqsell_tiled_kernel5()
387     slice_id     = row / SLICE_HEIGHT;  in matmult_seqsell_tiled_kernel4()
388     row_in_slice = row % SLICE_HEIGHT;  in matmult_seqsell_tiled_kernel4()
391 …r (i = sliidx[slice_id] + row_in_slice + SLICE_HEIGHT * threadIdx.y; i < sliidx[slice_id + 1]; i +…  in matmult_seqsell_tiled_kernel4()
411     slice_id     = row / SLICE_HEIGHT;  in matmult_seqsell_tiled_kernel3()
412     row_in_slice = row % SLICE_HEIGHT;  in matmult_seqsell_tiled_kernel3()
415 …r (i = sliidx[slice_id] + row_in_slice + SLICE_HEIGHT * threadIdx.y; i < sliidx[slice_id + 1]; i +…  in matmult_seqsell_tiled_kernel3()
433     slice_id     = row / SLICE_HEIGHT;  in matmult_seqsell_tiled_kernel2()
434     row_in_slice = row % SLICE_HEIGHT;  in matmult_seqsell_tiled_kernel2()
437 …r (i = sliidx[slice_id] + row_in_slice + SLICE_HEIGHT * threadIdx.y; i < sliidx[slice_id + 1]; i +…  in matmult_seqsell_tiled_kernel2()
453     slice_id     = row / SLICE_HEIGHT;  in matmultadd_seqsell_tiled_kernel6()
454     row_in_slice = row % SLICE_HEIGHT;  in matmultadd_seqsell_tiled_kernel6()
457 …r (i = sliidx[slice_id] + row_in_slice + SLICE_HEIGHT * threadIdx.y; i < sliidx[slice_id + 1]; i +…  in matmultadd_seqsell_tiled_kernel6()
481     slice_id     = row / SLICE_HEIGHT;  in matmultadd_seqsell_tiled_kernel5()
482     row_in_slice = row % SLICE_HEIGHT;  in matmultadd_seqsell_tiled_kernel5()
485 …r (i = sliidx[slice_id] + row_in_slice + SLICE_HEIGHT * threadIdx.y; i < sliidx[slice_id + 1]; i +…  in matmultadd_seqsell_tiled_kernel5()
507     slice_id     = row / SLICE_HEIGHT;  in matmultadd_seqsell_tiled_kernel4()
508     row_in_slice = row % SLICE_HEIGHT;  in matmultadd_seqsell_tiled_kernel4()
511 …r (i = sliidx[slice_id] + row_in_slice + SLICE_HEIGHT * threadIdx.y; i < sliidx[slice_id + 1]; i +…  in matmultadd_seqsell_tiled_kernel4()
531     slice_id     = row / SLICE_HEIGHT;  in matmultadd_seqsell_tiled_kernel3()
532     row_in_slice = row % SLICE_HEIGHT;  in matmultadd_seqsell_tiled_kernel3()
535 …r (i = sliidx[slice_id] + row_in_slice + SLICE_HEIGHT * threadIdx.y; i < sliidx[slice_id + 1]; i +…  in matmultadd_seqsell_tiled_kernel3()
553     slice_id     = row / SLICE_HEIGHT;  in matmultadd_seqsell_tiled_kernel2()
554     row_in_slice = row % SLICE_HEIGHT;  in matmultadd_seqsell_tiled_kernel2()
557 …r (i = sliidx[slice_id] + row_in_slice + SLICE_HEIGHT * threadIdx.y; i < sliidx[slice_id + 1]; i +…  in matmultadd_seqsell_tiled_kernel2()
585 …>kernelchoice >= 2 && cudastruct->kernelchoice <= 6 && sliceheight != SLICE_HEIGHT), PETSC_COMM_SE…  in MatMult_SeqSELLCUDA()
726 …>kernelchoice >= 2 && cudastruct->kernelchoice <= 6 && sliceheight != SLICE_HEIGHT), PETSC_COMM_SE…  in MatMultAdd_SeqSELLCUDA()