Lines Matching refs:y
78 …cInt *acolidx, const MatScalar *aval, const PetscInt *sliidx, const PetscScalar *x, PetscScalar *y) in matmult_seqsell_basic_kernel() argument
89 y[row] = sum; in matmult_seqsell_basic_kernel()
93 …atScalar *aval, const PetscInt *sliidx, const PetscScalar *x, const PetscScalar *y, PetscScalar *z) in matmultadd_seqsell_basic_kernel() argument
104 z[row] = y[row] + sum; in matmultadd_seqsell_basic_kernel()
112 …cInt *acolidx, const MatScalar *aval, const PetscInt *sliidx, const PetscScalar *x, PetscScalar *y) in matmult_seqsell_tiled_kernel9() argument
116 int tid = threadIdx.x + threadIdx.y * WARP_SIZE; in matmult_seqsell_tiled_kernel9()
124 …for (i = sliidx[slice_id] + threadIdx.x + WARP_SIZE * threadIdx.y; i < sliidx[slice_id + 1]; i += … in matmult_seqsell_tiled_kernel9()
129 if (threadIdx.x < sliceheight) shared[threadIdx.x][threadIdx.y] = t; in matmult_seqsell_tiled_kernel9()
136 if (row < nrows && threadIdx.y == 0 && threadIdx.x < sliceheight) y[row] = shared[0][threadIdx.x]; in matmult_seqsell_tiled_kernel9()
141 …atScalar *aval, const PetscInt *sliidx, const PetscScalar *x, const PetscScalar *y, PetscScalar *z) in matmultadd_seqsell_tiled_kernel9() argument
145 int tid = threadIdx.x + threadIdx.y * WARP_SIZE; in matmultadd_seqsell_tiled_kernel9()
153 …for (i = sliidx[slice_id] + threadIdx.x + WARP_SIZE * threadIdx.y; i < sliidx[slice_id + 1]; i += … in matmultadd_seqsell_tiled_kernel9()
158 if (threadIdx.x < sliceheight) shared[threadIdx.x][threadIdx.y] = t; in matmultadd_seqsell_tiled_kernel9()
165 …if (row < nrows && threadIdx.y == 0 && threadIdx.x < sliceheight) z[row] = y[row] + shared[0][thre… in matmultadd_seqsell_tiled_kernel9()
174 int halfwarpid = threadIdx.y * 2 + threadIdx.x / (WARP_SIZE / 2); in segment_scan()
175 shared[threadIdx.x + threadIdx.y * WARP_SIZE] = 0; in segment_scan()
177 shared[threadIdx.x + threadIdx.y * WARP_SIZE] = *val; in segment_scan()
181 …if (halfwarpid < BLOCKY * 2 - i) *val += shared[threadIdx.x + threadIdx.y * WARP_SIZE + i * WARP_S… in segment_scan()
189 …cInt *acolidx, const MatScalar *aval, const PetscInt *sliidx, const PetscScalar *x, PetscScalar *y) in matmult_seqsell_tiled_kernel8() argument
197 …gid = gridDim.x * WARP_SIZE * BLOCKY * iter + blockIdx.x * BLOCKY * WARP_SIZE + threadIdx.y * WARP… in matmult_seqsell_tiled_kernel8()
198 if (gid < nrows) y[gid] = 0.0; in matmult_seqsell_tiled_kernel8()
204 gid = cid * BLOCKY * WARP_SIZE + threadIdx.y * WARP_SIZE + threadIdx.x; in matmult_seqsell_tiled_kernel8()
211 …if (threadIdx.x % (WARP_SIZE / 2) == 0) flag[threadIdx.y * 2 + threadIdx.x / (WARP_SIZE / 2)] = sl… in matmult_seqsell_tiled_kernel8()
216 if (row < nrows && gid < totalentries && write) atomAdd(y[row], t); in matmult_seqsell_tiled_kernel8()
222 int tid = threadIdx.x + threadIdx.y * WARP_SIZE, tidx = tid % BLOCKY, tidy = tid / BLOCKY; in matmult_seqsell_tiled_kernel8()
227 …if (threadIdx.x < sliceheight) shared[threadIdx.x * BLOCKY + threadIdx.y] = t; /* shared[threadIdx… in matmult_seqsell_tiled_kernel8()
234 …if (row < nrows && threadIdx.y == 0 && threadIdx.x < sliceheight) atomAdd(y[row], shared[threadIdx… in matmult_seqsell_tiled_kernel8()
244 …atScalar *aval, const PetscInt *sliidx, const PetscScalar *x, const PetscScalar *y, PetscScalar *z) in matmultadd_seqsell_tiled_kernel8() argument
252 …gid = gridDim.x * WARP_SIZE * BLOCKY * iter + blockIdx.x * BLOCKY * WARP_SIZE + threadIdx.y * WARP… in matmultadd_seqsell_tiled_kernel8()
253 if (gid < nrows) z[gid] = y[gid]; in matmultadd_seqsell_tiled_kernel8()
259 gid = cid * BLOCKY * WARP_SIZE + threadIdx.y * WARP_SIZE + threadIdx.x; in matmultadd_seqsell_tiled_kernel8()
266 …if (threadIdx.x % (WARP_SIZE / 2) == 0) flag[threadIdx.y * 2 + threadIdx.x / (WARP_SIZE / 2)] = sl… in matmultadd_seqsell_tiled_kernel8()
277 int tid = threadIdx.x + threadIdx.y * WARP_SIZE, tidx = tid % BLOCKY, tidy = tid / BLOCKY; in matmultadd_seqsell_tiled_kernel8()
282 …if (threadIdx.x < sliceheight) shared[threadIdx.x * BLOCKY + threadIdx.y] = t; /* shared[threadIdx… in matmultadd_seqsell_tiled_kernel8()
289 …if (row < nrows && threadIdx.y == 0 && threadIdx.x < sliceheight) atomAdd(z[row], shared[threadIdx… in matmultadd_seqsell_tiled_kernel8()
298 …cInt *acolidx, const MatScalar *aval, const PetscInt *sliidx, const PetscScalar *x, PetscScalar *y) in matmult_seqsell_tiled_kernel7() argument
301 slice_id = blockIdx.x * blockDim.y + threadIdx.y; in matmult_seqsell_tiled_kernel7()
309 if (row < nrows && threadIdx.x < sliceheight) y[row] = t; in matmult_seqsell_tiled_kernel7()
313 …atScalar *aval, const PetscInt *sliidx, const PetscScalar *x, const PetscScalar *y, PetscScalar *z) in matmultadd_seqsell_tiled_kernel7() argument
316 slice_id = blockIdx.x * blockDim.y + threadIdx.y; in matmultadd_seqsell_tiled_kernel7()
324 if (row < nrows && threadIdx.x < sliceheight) z[row] = y[row] + t; in matmultadd_seqsell_tiled_kernel7()
331 …cInt *acolidx, const MatScalar *aval, const PetscInt *sliidx, const PetscScalar *x, PetscScalar *y) in PETSC_PRAGMA_DIAGNOSTIC_IGNORED_END()
341 shared[threadIdx.y * blockDim.x + threadIdx.x] = 0.0; in PETSC_PRAGMA_DIAGNOSTIC_IGNORED_END()
342 …w_in_slice + sliceheight * threadIdx.y; i < sliidx[slice_id + 1]; i += sliceheight * blockDim.y) s… in PETSC_PRAGMA_DIAGNOSTIC_IGNORED_END()
344 …if (threadIdx.y < 16) shared[threadIdx.y * blockDim.x + threadIdx.x] += shared[(threadIdx.y + 16) … in PETSC_PRAGMA_DIAGNOSTIC_IGNORED_END()
346 …if (threadIdx.y < 8) shared[threadIdx.y * blockDim.x + threadIdx.x] += shared[(threadIdx.y + 8) * … in PETSC_PRAGMA_DIAGNOSTIC_IGNORED_END()
348 …if (threadIdx.y < 4) shared[threadIdx.y * blockDim.x + threadIdx.x] += shared[(threadIdx.y + 4) * … in PETSC_PRAGMA_DIAGNOSTIC_IGNORED_END()
350 …if (threadIdx.y < 2) shared[threadIdx.y * blockDim.x + threadIdx.x] += shared[(threadIdx.y + 2) * … in PETSC_PRAGMA_DIAGNOSTIC_IGNORED_END()
352 if (threadIdx.y < 1) { in PETSC_PRAGMA_DIAGNOSTIC_IGNORED_END()
354 y[row] = shared[threadIdx.x]; in PETSC_PRAGMA_DIAGNOSTIC_IGNORED_END()
359 …cInt *acolidx, const MatScalar *aval, const PetscInt *sliidx, const PetscScalar *x, PetscScalar *y) in matmult_seqsell_tiled_kernel5() argument
369 shared[threadIdx.y * blockDim.x + threadIdx.x] = 0.0; in matmult_seqsell_tiled_kernel5()
370 …w_in_slice + sliceheight * threadIdx.y; i < sliidx[slice_id + 1]; i += sliceheight * blockDim.y) s… in matmult_seqsell_tiled_kernel5()
372 …if (threadIdx.y < 8) shared[threadIdx.y * blockDim.x + threadIdx.x] += shared[(threadIdx.y + 8) * … in matmult_seqsell_tiled_kernel5()
374 …if (threadIdx.y < 4) shared[threadIdx.y * blockDim.x + threadIdx.x] += shared[(threadIdx.y + 4) * … in matmult_seqsell_tiled_kernel5()
376 …if (threadIdx.y < 2) shared[threadIdx.y * blockDim.x + threadIdx.x] += shared[(threadIdx.y + 2) * … in matmult_seqsell_tiled_kernel5()
378 if (threadIdx.y < 1) { in matmult_seqsell_tiled_kernel5()
380 y[row] = shared[threadIdx.x]; in matmult_seqsell_tiled_kernel5()
385 …cInt *acolidx, const MatScalar *aval, const PetscInt *sliidx, const PetscScalar *x, PetscScalar *y) in matmult_seqsell_tiled_kernel4() argument
395 shared[threadIdx.y * blockDim.x + threadIdx.x] = 0.0; in matmult_seqsell_tiled_kernel4()
396 …w_in_slice + sliceheight * threadIdx.y; i < sliidx[slice_id + 1]; i += sliceheight * blockDim.y) s… in matmult_seqsell_tiled_kernel4()
398 …if (threadIdx.y < 4) shared[threadIdx.y * blockDim.x + threadIdx.x] += shared[(threadIdx.y + 4) * … in matmult_seqsell_tiled_kernel4()
400 …if (threadIdx.y < 2) shared[threadIdx.y * blockDim.x + threadIdx.x] += shared[(threadIdx.y + 2) * … in matmult_seqsell_tiled_kernel4()
402 if (threadIdx.y < 1) { in matmult_seqsell_tiled_kernel4()
404 y[row] = shared[threadIdx.x]; in matmult_seqsell_tiled_kernel4()
409 …cInt *acolidx, const MatScalar *aval, const PetscInt *sliidx, const PetscScalar *x, PetscScalar *y) in matmult_seqsell_tiled_kernel3() argument
419 shared[threadIdx.y * blockDim.x + threadIdx.x] = 0.0; in matmult_seqsell_tiled_kernel3()
420 …w_in_slice + sliceheight * threadIdx.y; i < sliidx[slice_id + 1]; i += sliceheight * blockDim.y) s… in matmult_seqsell_tiled_kernel3()
422 …if (threadIdx.y < 2) shared[threadIdx.y * blockDim.x + threadIdx.x] += shared[(threadIdx.y + 2) * … in matmult_seqsell_tiled_kernel3()
424 if (threadIdx.y < 1) { in matmult_seqsell_tiled_kernel3()
426 y[row] = shared[threadIdx.x]; in matmult_seqsell_tiled_kernel3()
431 …cInt *acolidx, const MatScalar *aval, const PetscInt *sliidx, const PetscScalar *x, PetscScalar *y) in matmult_seqsell_tiled_kernel2() argument
441 shared[threadIdx.y * blockDim.x + threadIdx.x] = 0.0; in matmult_seqsell_tiled_kernel2()
442 …w_in_slice + sliceheight * threadIdx.y; i < sliidx[slice_id + 1]; i += sliceheight * blockDim.y) s… in matmult_seqsell_tiled_kernel2()
444 if (threadIdx.y < 1) { in matmult_seqsell_tiled_kernel2()
446 y[row] = shared[threadIdx.x]; in matmult_seqsell_tiled_kernel2()
451 …atScalar *aval, const PetscInt *sliidx, const PetscScalar *x, const PetscScalar *y, PetscScalar *z) in matmultadd_seqsell_tiled_kernel6() argument
461 shared[threadIdx.y * blockDim.x + threadIdx.x] = 0.0; in matmultadd_seqsell_tiled_kernel6()
462 …w_in_slice + sliceheight * threadIdx.y; i < sliidx[slice_id + 1]; i += sliceheight * blockDim.y) s… in matmultadd_seqsell_tiled_kernel6()
464 …if (threadIdx.y < 16) shared[threadIdx.y * blockDim.x + threadIdx.x] += shared[(threadIdx.y + 16) … in matmultadd_seqsell_tiled_kernel6()
466 …if (threadIdx.y < 8) shared[threadIdx.y * blockDim.x + threadIdx.x] += shared[(threadIdx.y + 8) * … in matmultadd_seqsell_tiled_kernel6()
468 …if (threadIdx.y < 4) shared[threadIdx.y * blockDim.x + threadIdx.x] += shared[(threadIdx.y + 4) * … in matmultadd_seqsell_tiled_kernel6()
470 …if (threadIdx.y < 2) shared[threadIdx.y * blockDim.x + threadIdx.x] += shared[(threadIdx.y + 2) * … in matmultadd_seqsell_tiled_kernel6()
472 if (threadIdx.y < 1) { in matmultadd_seqsell_tiled_kernel6()
474 z[row] = y[row] + shared[threadIdx.x]; in matmultadd_seqsell_tiled_kernel6()
479 …atScalar *aval, const PetscInt *sliidx, const PetscScalar *x, const PetscScalar *y, PetscScalar *z) in matmultadd_seqsell_tiled_kernel5() argument
489 shared[threadIdx.y * blockDim.x + threadIdx.x] = 0.0; in matmultadd_seqsell_tiled_kernel5()
490 …w_in_slice + sliceheight * threadIdx.y; i < sliidx[slice_id + 1]; i += sliceheight * blockDim.y) s… in matmultadd_seqsell_tiled_kernel5()
492 …if (threadIdx.y < 8) shared[threadIdx.y * blockDim.x + threadIdx.x] += shared[(threadIdx.y + 8) * … in matmultadd_seqsell_tiled_kernel5()
494 …if (threadIdx.y < 4) shared[threadIdx.y * blockDim.x + threadIdx.x] += shared[(threadIdx.y + 4) * … in matmultadd_seqsell_tiled_kernel5()
496 …if (threadIdx.y < 2) shared[threadIdx.y * blockDim.x + threadIdx.x] += shared[(threadIdx.y + 2) * … in matmultadd_seqsell_tiled_kernel5()
498 if (threadIdx.y < 1) { in matmultadd_seqsell_tiled_kernel5()
500 z[row] = y[row] + shared[threadIdx.x]; in matmultadd_seqsell_tiled_kernel5()
505 …atScalar *aval, const PetscInt *sliidx, const PetscScalar *x, const PetscScalar *y, PetscScalar *z) in matmultadd_seqsell_tiled_kernel4() argument
515 shared[threadIdx.y * blockDim.x + threadIdx.x] = 0.0; in matmultadd_seqsell_tiled_kernel4()
516 …w_in_slice + sliceheight * threadIdx.y; i < sliidx[slice_id + 1]; i += sliceheight * blockDim.y) s… in matmultadd_seqsell_tiled_kernel4()
518 …if (threadIdx.y < 4) shared[threadIdx.y * blockDim.x + threadIdx.x] += shared[(threadIdx.y + 4) * … in matmultadd_seqsell_tiled_kernel4()
520 …if (threadIdx.y < 2) shared[threadIdx.y * blockDim.x + threadIdx.x] += shared[(threadIdx.y + 2) * … in matmultadd_seqsell_tiled_kernel4()
522 if (threadIdx.y < 1) { in matmultadd_seqsell_tiled_kernel4()
524 z[row] = y[row] + shared[threadIdx.x]; in matmultadd_seqsell_tiled_kernel4()
529 …atScalar *aval, const PetscInt *sliidx, const PetscScalar *x, const PetscScalar *y, PetscScalar *z) in matmultadd_seqsell_tiled_kernel3() argument
539 shared[threadIdx.y * blockDim.x + threadIdx.x] = 0.0; in matmultadd_seqsell_tiled_kernel3()
540 …w_in_slice + sliceheight * threadIdx.y; i < sliidx[slice_id + 1]; i += sliceheight * blockDim.y) s… in matmultadd_seqsell_tiled_kernel3()
542 …if (threadIdx.y < 2) shared[threadIdx.y * blockDim.x + threadIdx.x] += shared[(threadIdx.y + 2) * … in matmultadd_seqsell_tiled_kernel3()
544 if (threadIdx.y < 1) { in matmultadd_seqsell_tiled_kernel3()
546 z[row] = y[row] + shared[threadIdx.x]; in matmultadd_seqsell_tiled_kernel3()
551 …atScalar *aval, const PetscInt *sliidx, const PetscScalar *x, const PetscScalar *y, PetscScalar *z) in matmultadd_seqsell_tiled_kernel2() argument
561 shared[threadIdx.y * blockDim.x + threadIdx.x] = 0.0; in matmultadd_seqsell_tiled_kernel2()
562 …w_in_slice + sliceheight * threadIdx.y; i < sliidx[slice_id + 1]; i += sliceheight * blockDim.y) s… in matmultadd_seqsell_tiled_kernel2()
564 if (threadIdx.y < 1) { in matmultadd_seqsell_tiled_kernel2()
566 z[row] = y[row] + shared[threadIdx.x]; in matmultadd_seqsell_tiled_kernel2()
575 PetscScalar *y; in MatMult_SeqSELLHIP() local
598 PetscCall(VecHIPGetArrayWrite(yy, &y)); in MatMult_SeqSELLHIP()
606 …tiled_kernel9<2><<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
608 …tiled_kernel9<4><<<nblocks, dim3(WARP_SIZE, 4)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
610 …tiled_kernel9<8><<<nblocks, dim3(WARP_SIZE, 8)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
612 …led_kernel9<16><<<nblocks, dim3(WARP_SIZE, 16)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
614 …tiled_kernel9<2><<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
620 …ll_tiled_kernel7<<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
622 …ll_tiled_kernel7<<<nblocks, dim3(WARP_SIZE, 4)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
624 …ll_tiled_kernel7<<<nblocks, dim3(WARP_SIZE, 8)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
626 …l_tiled_kernel7<<<nblocks, dim3(WARP_SIZE, 16)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
629 …ll_tiled_kernel7<<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
635 …tmult_seqsell_tiled_kernel6<<<nblocks, block32>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
639 …tmult_seqsell_tiled_kernel5<<<nblocks, block16>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
643 …atmult_seqsell_tiled_kernel4<<<nblocks, block8>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
647 …atmult_seqsell_tiled_kernel3<<<nblocks, block4>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
651 …atmult_seqsell_tiled_kernel2<<<nblocks, block2>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
655 …mult_seqsell_basic_kernel<<<nblocks, blocksize>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
668 …, 2)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
670 …, 4)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
672 …, 8)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
674 … 16)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
676 …, 2)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
683 …ll_tiled_kernel7<<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
686 …tiled_kernel9<2><<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
690 …tiled_kernel9<8><<<nblocks, dim3(WARP_SIZE, 8)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
693 …led_kernel9<16><<<nblocks, dim3(WARP_SIZE, 16)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y); in MatMult_SeqSELLHIP()
703 PetscCall(VecHIPRestoreArrayWrite(yy, &y)); in MatMult_SeqSELLHIP()
713 const PetscScalar *y, *x; in MatMultAdd_SeqSELLHIP() local
732 PetscCall(VecHIPGetArrayRead(yy, &y)); in MatMultAdd_SeqSELLHIP()
741 …ed_kernel9<2><<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
743 …ed_kernel9<4><<<nblocks, dim3(WARP_SIZE, 4)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
745 …ed_kernel9<8><<<nblocks, dim3(WARP_SIZE, 8)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
747 …_kernel9<16><<<nblocks, dim3(WARP_SIZE, 16)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
749 …ed_kernel9<2><<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
760 …)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
762 …)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
764 …)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
766 …)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
768 …)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
774 …tiled_kernel7<<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
776 …tiled_kernel7<<<nblocks, dim3(WARP_SIZE, 4)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
778 …tiled_kernel7<<<nblocks, dim3(WARP_SIZE, 8)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
780 …iled_kernel7<<<nblocks, dim3(WARP_SIZE, 16)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
783 …tiled_kernel7<<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
789 …dd_seqsell_tiled_kernel6<<<nblocks, block32>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
793 …dd_seqsell_tiled_kernel5<<<nblocks, block16>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
797 …add_seqsell_tiled_kernel4<<<nblocks, block8>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
801 …add_seqsell_tiled_kernel3<<<nblocks, block4>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
805 …add_seqsell_tiled_kernel2<<<nblocks, block2>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
809 …d_seqsell_basic_kernel<<<nblocks, blocksize>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
822 …)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
824 …)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
826 …)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
828 …)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
830 …)>>>(nrows, sliceheight, chunksperblock, nchunks, chunk_slice_map, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
837 …tiled_kernel7<<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
840 …ed_kernel9<2><<<nblocks, dim3(WARP_SIZE, 2)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
844 …ed_kernel9<8><<<nblocks, dim3(WARP_SIZE, 8)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
847 …_kernel9<16><<<nblocks, dim3(WARP_SIZE, 16)>>>(nrows, sliceheight, acolidx, aval, sliidx, x, y, z); in MatMultAdd_SeqSELLHIP()
857 PetscCall(VecHIPRestoreArrayRead(yy, &y)); in MatMultAdd_SeqSELLHIP()