Lines Matching refs:gid
188 PetscInt gid, row, start_slice, cid; in matmult_seqsell_tiled_kernel8() local
193 … gid = gridDim.x * 32 * BLOCKY * iter + blockIdx.x * BLOCKY * 32 + threadIdx.y * 32 + threadIdx.x; in matmult_seqsell_tiled_kernel8()
194 if (gid < nrows) y[gid] = 0.0; in matmult_seqsell_tiled_kernel8()
200 gid = cid * BLOCKY * 32 + threadIdx.y * 32 + threadIdx.x; in matmult_seqsell_tiled_kernel8()
206 while (gid < totalentries && gid >= sliidx[slice_id + 1]) slice_id++; in matmult_seqsell_tiled_kernel8()
209 if (row < nrows && gid < totalentries) t = aval[gid] * x[acolidx[gid]]; in matmult_seqsell_tiled_kernel8()
212 if (row < nrows && gid < totalentries && write) atomAdd(y[row], t); in matmult_seqsell_tiled_kernel8()
216 if (row < nrows) t += aval[gid] * x[acolidx[gid]]; in matmult_seqsell_tiled_kernel8()
243 PetscInt gid, row, start_slice, cid; in matmultadd_seqsell_tiled_kernel8() local
248 … gid = gridDim.x * 32 * BLOCKY * iter + blockIdx.x * BLOCKY * 32 + threadIdx.y * 32 + threadIdx.x; in matmultadd_seqsell_tiled_kernel8()
249 if (gid < nrows) z[gid] = y[gid]; in matmultadd_seqsell_tiled_kernel8()
255 gid = cid * BLOCKY * 32 + threadIdx.y * 32 + threadIdx.x; in matmultadd_seqsell_tiled_kernel8()
261 while (gid < totalentries && gid >= sliidx[slice_id + 1]) slice_id++; in matmultadd_seqsell_tiled_kernel8()
264 if (row < nrows && gid < totalentries) t = aval[gid] * x[acolidx[gid]]; in matmultadd_seqsell_tiled_kernel8()
267 if (row < nrows && gid < totalentries && write) atomAdd(z[row], t); in matmultadd_seqsell_tiled_kernel8()
271 if (row < nrows) t += aval[gid] * x[acolidx[gid]]; in matmultadd_seqsell_tiled_kernel8()