Lines Matching refs:hipstruct
25 static PetscErrorCode MatSeqSELLHIP_Destroy(Mat_SeqSELLHIP **hipstruct) in MatSeqSELLHIP_Destroy() argument
28 if (*hipstruct) { in MatSeqSELLHIP_Destroy()
29 if ((*hipstruct)->colidx) PetscCallHIP(hipFree((*hipstruct)->colidx)); in MatSeqSELLHIP_Destroy()
30 if ((*hipstruct)->val) PetscCallHIP(hipFree((*hipstruct)->val)); in MatSeqSELLHIP_Destroy()
31 if ((*hipstruct)->sliidx) PetscCallHIP(hipFree((*hipstruct)->sliidx)); in MatSeqSELLHIP_Destroy()
32 if ((*hipstruct)->chunk_slice_map) PetscCallHIP(hipFree((*hipstruct)->chunk_slice_map)); in MatSeqSELLHIP_Destroy()
33 PetscCall(PetscFree(*hipstruct)); in MatSeqSELLHIP_Destroy()
40 Mat_SeqSELLHIP *hipstruct = (Mat_SeqSELLHIP *)A->spptr; in MatSeqSELLHIPCopyToGPU() local
46 if (A->assembled && A->nonzerostate == hipstruct->nonzerostate) { in MatSeqSELLHIPCopyToGPU()
48 …PetscCallHIP(hipMemcpy(hipstruct->val, a->val, a->sliidx[a->totalslices] * sizeof(MatScalar), hipM… in MatSeqSELLHIPCopyToGPU()
51 if (hipstruct->colidx) PetscCallHIP(hipFree(hipstruct->colidx)); in MatSeqSELLHIPCopyToGPU()
52 if (hipstruct->val) PetscCallHIP(hipFree(hipstruct->val)); in MatSeqSELLHIPCopyToGPU()
53 if (hipstruct->sliidx) PetscCallHIP(hipFree(hipstruct->sliidx)); in MatSeqSELLHIPCopyToGPU()
54 if (hipstruct->chunk_slice_map) PetscCallHIP(hipFree(hipstruct->chunk_slice_map)); in MatSeqSELLHIPCopyToGPU()
55 hipstruct->maxallocmat = a->maxallocmat; in MatSeqSELLHIPCopyToGPU()
56 hipstruct->totalentries = a->sliidx[a->totalslices]; in MatSeqSELLHIPCopyToGPU()
57 hipstruct->totalslices = a->totalslices; in MatSeqSELLHIPCopyToGPU()
58 hipstruct->totalchunks = a->totalchunks; in MatSeqSELLHIPCopyToGPU()
59 … PetscCallHIP(hipMalloc((void **)&hipstruct->colidx, a->maxallocmat * sizeof(*hipstruct->colidx))); in MatSeqSELLHIPCopyToGPU()
60 PetscCallHIP(hipMalloc((void **)&hipstruct->val, a->maxallocmat * sizeof(*hipstruct->val))); in MatSeqSELLHIPCopyToGPU()
62 …PetscCallHIP(hipMemcpy(hipstruct->colidx, a->colidx, a->sliidx[a->totalslices] * sizeof(*a->colidx… in MatSeqSELLHIPCopyToGPU()
63 …PetscCallHIP(hipMemcpy(hipstruct->val, a->val, a->sliidx[a->totalslices] * sizeof(*a->val), hipMem… in MatSeqSELLHIPCopyToGPU()
65 …PetscCallHIP(hipMalloc((void **)&hipstruct->sliidx, (a->totalslices + 1) * sizeof(*hipstruct->slii… in MatSeqSELLHIPCopyToGPU()
66 …PetscCallHIP(hipMemcpy(hipstruct->sliidx, a->sliidx, (a->totalslices + 1) * sizeof(*a->sliidx), hi… in MatSeqSELLHIPCopyToGPU()
67 …PetscCallHIP(hipMalloc((void **)&hipstruct->chunk_slice_map, a->totalchunks * sizeof(*hipstruct->c… in MatSeqSELLHIPCopyToGPU()
68 …PetscCallHIP(hipMemcpy(hipstruct->chunk_slice_map, a->chunk_slice_map, a->totalchunks * sizeof(*a-… in MatSeqSELLHIPCopyToGPU()
574 Mat_SeqSELLHIP *hipstruct = (Mat_SeqSELLHIP *)A->spptr; in MatMult_SeqSELLHIP() local
590 …PetscCheck(!(hipstruct->kernelchoice >= 2 && hipstruct->kernelchoice <= 6 && sliceheight > 32), PE… in MatMult_SeqSELLHIP()
593 aval = hipstruct->val; in MatMult_SeqSELLHIP()
594 acolidx = hipstruct->colidx; in MatMult_SeqSELLHIP()
595 sliidx = hipstruct->sliidx; in MatMult_SeqSELLHIP()
601 switch (hipstruct->kernelchoice) { in MatMult_SeqSELLHIP()
605 if (hipstruct->blocky == 2) { in MatMult_SeqSELLHIP()
607 } else if (hipstruct->blocky == 4) { in MatMult_SeqSELLHIP()
609 } else if (hipstruct->blocky == 8) { in MatMult_SeqSELLHIP()
611 } else if (hipstruct->blocky == 16) { in MatMult_SeqSELLHIP()
618 nblocks = 1 + (nrows - 1) / (hipstruct->blocky * sliceheight); in MatMult_SeqSELLHIP()
619 if (hipstruct->blocky == 2) { in MatMult_SeqSELLHIP()
621 } else if (hipstruct->blocky == 4) { in MatMult_SeqSELLHIP()
623 } else if (hipstruct->blocky == 8) { in MatMult_SeqSELLHIP()
625 } else if (hipstruct->blocky == 16) { in MatMult_SeqSELLHIP()
663 nchunks = hipstruct->totalchunks; in MatMult_SeqSELLHIP()
664 …chunksperblock = hipstruct->chunksperblock ? hipstruct->chunksperblock : 1 + (hipstruct->totalent… in MatMult_SeqSELLHIP()
666 chunk_slice_map = hipstruct->chunk_slice_map; in MatMult_SeqSELLHIP()
699 …P, "unsupported kernel choice %" PetscInt_FMT " for MatMult_SeqSELLHIP.", hipstruct->kernelchoice); in MatMult_SeqSELLHIP()
711 Mat_SeqSELLHIP *hipstruct = (Mat_SeqSELLHIP *)A->spptr; in MatMultAdd_SeqSELLHIP() local
715 MatScalar *aval = hipstruct->val; in MatMultAdd_SeqSELLHIP()
716 PetscInt *acolidx = hipstruct->colidx; in MatMultAdd_SeqSELLHIP()
717 PetscInt *sliidx = hipstruct->sliidx; in MatMultAdd_SeqSELLHIP()
721 PetscInt blocky = hipstruct->blocky; in MatMultAdd_SeqSELLHIP()
726 …PetscCheck(!(hipstruct->kernelchoice >= 2 && hipstruct->kernelchoice <= 6 && sliceheight != sliceh… in MatMultAdd_SeqSELLHIP()
736 switch (hipstruct->kernelchoice) { in MatMultAdd_SeqSELLHIP()
754 nchunks = hipstruct->totalchunks; in MatMultAdd_SeqSELLHIP()
756 …chunksperblock = hipstruct->chunksperblock ? hipstruct->chunksperblock : 1 + (hipstruct->totalent… in MatMultAdd_SeqSELLHIP()
758 chunk_slice_map = hipstruct->chunk_slice_map; in MatMultAdd_SeqSELLHIP()
816 nchunks = hipstruct->totalchunks; in MatMultAdd_SeqSELLHIP()
818 …chunksperblock = hipstruct->chunksperblock ? hipstruct->chunksperblock : 1 + (hipstruct->totalent… in MatMultAdd_SeqSELLHIP()
820 chunk_slice_map = hipstruct->chunk_slice_map; in MatMultAdd_SeqSELLHIP()
853 …P, "unsupported kernel choice %" PetscInt_FMT " for MatMult_SeqSELLHIP.", hipstruct->kernelchoice); in MatMultAdd_SeqSELLHIP()
868 Mat_SeqSELLHIP *hipstruct = (Mat_SeqSELLHIP *)A->spptr; in MatSetFromOptions_SeqSELLHIP() local
877 hipstruct->blocky = blocky; in MatSetFromOptions_SeqSELLHIP()
882 hipstruct->kernelchoice = kernel; in MatSetFromOptions_SeqSELLHIP()
883 …ll(PetscOptionsGetInt(NULL, NULL, "-mat_sell_spmv_hip_chunksperblock", &hipstruct->chunksperblock,… in MatSetFromOptions_SeqSELLHIP()
919 Mat_SeqSELLHIP *hipstruct = (Mat_SeqSELLHIP *)A->spptr; in MatZeroEntries_SeqSELLHIP() local
920 if (hipstruct->val) { in MatZeroEntries_SeqSELLHIP()
922 … PetscCallHIP(hipMemset(hipstruct->val, 0, a->sliidx[a->totalslices] * sizeof(*hipstruct->val))); in MatZeroEntries_SeqSELLHIP()
950 Mat_SeqSELLHIP *hipstruct; in MatConvert_SeqSELL_SeqSELLHIP() local
958 PetscCall(PetscNew(&hipstruct)); in MatConvert_SeqSELL_SeqSELLHIP()
959 B->spptr = hipstruct; in MatConvert_SeqSELL_SeqSELLHIP()