Home
last modified time | relevance | path

Searched refs:sliceheight (Results 1 – 6 of 6) sorted by relevance

/petsc/src/mat/impls/sell/seq/seqhip/
H A Dsellhip.hip.cxx78 static __global__ void matmult_seqsell_basic_kernel(PetscInt nrows, PetscInt sliceheight, const Pet… in matmult_seqsell_basic_kernel() argument
85 slice_id = row / sliceheight; in matmult_seqsell_basic_kernel()
86 row_in_slice = row % sliceheight; in matmult_seqsell_basic_kernel()
88 …for (i = sliidx[slice_id] + row_in_slice; i < sliidx[slice_id + 1]; i += sliceheight) sum += aval[… in matmult_seqsell_basic_kernel()
93 static __global__ void matmultadd_seqsell_basic_kernel(PetscInt nrows, PetscInt sliceheight, const … in matmultadd_seqsell_basic_kernel() argument
100 slice_id = row / sliceheight; in matmultadd_seqsell_basic_kernel()
101 row_in_slice = row % sliceheight; in matmultadd_seqsell_basic_kernel()
103 …for (i = sliidx[slice_id] + row_in_slice; i < sliidx[slice_id + 1]; i += sliceheight) sum += aval[… in matmultadd_seqsell_basic_kernel()
112 __global__ void matmult_seqsell_tiled_kernel9(PetscInt nrows, PetscInt sliceheight, const PetscInt … in matmult_seqsell_tiled_kernel9() argument
122 row = slice_id * sliceheight + threadIdx.x % sliceheight; in matmult_seqsell_tiled_kernel9()
[all …]
/petsc/src/mat/impls/sell/seq/
H A Dsell.h47 PetscInt sliceheight; /* slice height */ \
134 if (*(cp + a->sliceheight * t) > col) high = t; \
138 if (*(cp + a->sliceheight * _i) > col) break; \
139 if (*(cp + a->sliceheight * _i) == col) { \
140 if (addv == ADD_VALUES) *(vp + a->sliceheight * _i) += value; \
141 else *(vp + a->sliceheight * _i) = value; \
148 …s) && a->rlen[row] >= (a->sliidx[row / a->sliceheight + 1] - a->sliidx[row / a->sliceheight]) / a-…
150 if (a->maxallocmat < a->sliidx[a->totalslices] + a->sliceheight) { \
152 PetscInt new_size = a->maxallocmat + a->sliceheight, *new_colidx; \
158 PetscCall(PetscArraycpy(new_val, a->val, a->sliidx[row / a->sliceheight + 1])); \
[all …]
H A Dsell.c120 if (!b->sliceheight) { /* not set yet */ in MatSeqSELLSetPreallocation_SeqSELL()
122 b->sliceheight = 16; in MatSeqSELLSetPreallocation_SeqSELL()
124 b->sliceheight = 8; in MatSeqSELLSetPreallocation_SeqSELL()
127 totalslices = PetscCeilInt(B->rmap->n, b->sliceheight); in MatSeqSELLSetPreallocation_SeqSELL()
130 …if (B->rmap->n % b->sliceheight) PetscCall(PetscInfo(B, "Padding rows to the SEQSELL matrix becaus… in MatSeqSELLSetPreallocation_SeqSELL()
141 while (b->sliceheight * maxallocrow % DEVICE_MEM_ALIGN) maxallocrow++; in MatSeqSELLSetPreallocation_SeqSELL()
143 for (i = 0; i <= totalslices; i++) b->sliidx[i] = b->sliceheight * i * maxallocrow; in MatSeqSELLSetPreallocation_SeqSELL()
146 PetscInt mul = DEVICE_MEM_ALIGN / b->sliceheight; in MatSeqSELLSetPreallocation_SeqSELL()
152 …for (j = 0; j < b->sliceheight; j++) b->sliidx[i] = PetscMax(b->sliidx[i], rlen[b->sliceheight * (… in MatSeqSELLSetPreallocation_SeqSELL()
160 PetscCall(PetscIntSumError(b->sliidx[i - 1], b->sliceheight * b->sliidx[i], &b->sliidx[i])); in MatSeqSELLSetPreallocation_SeqSELL()
[all …]
/petsc/src/mat/impls/sell/seq/seqcuda/
H A Dsellcuda.cu75 static __global__ void matmult_seqsell_basic_kernel(PetscInt nrows, PetscInt sliceheight, const Pet… in matmult_seqsell_basic_kernel() argument
82 slice_id = row / sliceheight; in matmult_seqsell_basic_kernel()
83 row_in_slice = row % sliceheight; in matmult_seqsell_basic_kernel()
85 …for (i = sliidx[slice_id] + row_in_slice; i < sliidx[slice_id + 1]; i += sliceheight) sum += aval[… in matmult_seqsell_basic_kernel()
90 static __global__ void matmultadd_seqsell_basic_kernel(PetscInt nrows, PetscInt sliceheight, const … in matmultadd_seqsell_basic_kernel() argument
97 slice_id = row / sliceheight; in matmultadd_seqsell_basic_kernel()
98 row_in_slice = row % sliceheight; in matmultadd_seqsell_basic_kernel()
100 …for (i = sliidx[slice_id] + row_in_slice; i < sliidx[slice_id + 1]; i += sliceheight) sum += aval[… in matmultadd_seqsell_basic_kernel()
108 __global__ void matmult_seqsell_tiled_kernel9(PetscInt nrows, PetscInt sliceheight, const PetscInt … in matmult_seqsell_tiled_kernel9() argument
118 row = slice_id * sliceheight + threadIdx.x % sliceheight; in matmult_seqsell_tiled_kernel9()
[all …]
/petsc/src/mat/impls/sell/mpi/
H A Dmmsell.c56 totalslices = PetscCeilInt(B->rmap->n, Bsell->sliceheight); in MatDisAssemble_MPISELL()
58 …= Bsell->sliidx[i], row = 0; j < Bsell->sliidx[i + 1]; j++, row = (row + 1) % Bsell->sliceheight) { in MatDisAssemble_MPISELL()
59 …nonzero = (PetscBool)((j - Bsell->sliidx[i]) / Bsell->sliceheight < Bsell->rlen[Bsell->sliceheight… in MatDisAssemble_MPISELL()
60 …if (isnonzero) PetscCall(MatSetValue(Bnew, Bsell->sliceheight * i + row, sell->garray[Bsell->colid… in MatDisAssemble_MPISELL()
90 totalslices = PetscCeilInt(sell->B->rmap->n, B->sliceheight); in MatSetUpMultiply_MPISELL()
98 …nonzero = (PetscBool)((j - B->sliidx[i]) / B->sliceheight < B->rlen[i * B->sliceheight + j % B->sl… in MatSetUpMultiply_MPISELL()
127 …nonzero = (PetscBool)((j - B->sliidx[i]) / B->sliceheight < B->rlen[i * B->sliceheight + j % B->sl… in MatSetUpMultiply_MPISELL()
145 …nonzero = (PetscBool)((j - B->sliidx[i]) / B->sliceheight < B->rlen[i * B->sliceheight + j % B->sl… in MatSetUpMultiply_MPISELL()
166 …nonzero = (PetscBool)((j - B->sliidx[i]) / B->sliceheight < B->rlen[i * B->sliceheight + j % B->sl… in MatSetUpMultiply_MPISELL()
H A Dmpisell.c69 if (cp1[sliceheight * t] > col) high1 = t; \
73 if (cp1[sliceheight * _i] > col) break; \
74 if (cp1[sliceheight * _i] == col) { \
75 if (addv == ADD_VALUES) vp1[sliceheight * _i] += value; \
76 else vp1[sliceheight * _i] = value; \
92 …MatSeqXSELLReallocateSELL(A, am, 1, nrow1, a->sliidx, a->sliceheight, row / sliceheight, row, col,…
95 cp1[sliceheight * (ii + 1)] = cp1[sliceheight * ii]; \
96 vp1[sliceheight * (ii + 1)] = vp1[sliceheight * ii]; \
98 cp1[sliceheight * _i] = col; \
99 vp1[sliceheight * _i] = value; \
[all …]