diff --git a/.gitignore b/.gitignore index 3fcd9db..2fe5eed 100644 --- a/.gitignore +++ b/.gitignore @@ -1,4 +1,4 @@ build/ .cache/ .vscode/ -URDepends/ \ No newline at end of file +URD*/ \ No newline at end of file diff --git a/TVALGPU/src/mat_vec_mul.h b/TVALGPU/src/mat_vec_mul.h index 0903554..b9cdde1 100644 --- a/TVALGPU/src/mat_vec_mul.h +++ b/TVALGPU/src/mat_vec_mul.h @@ -41,8 +41,8 @@ struct sparse_mm { cusparseIndexBase_t idxBase) */ - - HANDLE_ERROR(cusparseScsr2csc(cs_handle, + unsigned long tbufferSize = 0; + cusparseCsr2cscEx2_bufferSize(cs_handle, A.dim_x, // m A.dim_y, // n A.nnz, // nnz @@ -50,10 +50,42 @@ struct sparse_mm { A_csc.ptr(), // csrRowPtr A_csc.ind(), // csrColInd A_csr.val(), // cscVal +A_csr.ptr(), // cscColPtr A_csr.ind(), // cscRowInd +CUDA_R_32F, + CUSPARSE_ACTION_NUMERIC, // copyValues + CUSPARSE_INDEX_BASE_ZERO, + CUSPARSE_CSR2CSC_ALG1, + &tbufferSize); + void * buffer =NULL; + cudaMalloc(&buffer,tbufferSize); + cusparseCsr2cscEx2(cs_handle, + A.dim_x, // m + A.dim_y, // n + A.nnz, // nnz + A_csc.val(), // csrVal + A_csc.ptr(), // csrRowPtr + A_csc.ind(), // csrColInd + A_csr.val(), // cscVal A_csr.ptr(), // cscColPtr + A_csr.ind(), // cscRowInd + CUDA_R_32F, CUSPARSE_ACTION_NUMERIC, // copyValues - CUSPARSE_INDEX_BASE_ZERO)); // idxBase + CUSPARSE_INDEX_BASE_ZERO, + CUSPARSE_CSR2CSC_ALG1,buffer); + cudaFree(buffer); + // HANDLE_ERROR(cusparseScsr2csc(cs_handle, + // A.dim_x, // m + // A.dim_y, // n + // A.nnz, // nnz + // A_csc.val(), // csrVal + // A_csc.ptr(), // csrRowPtr + // A_csc.ind(), // csrColInd + // A_csr.val(), // cscVal + // A_csr.ind(), // cscRowInd + // A_csr.ptr(), // cscColPtr + // CUSPARSE_ACTION_NUMERIC, // copyValues + // CUSPARSE_INDEX_BASE_ZERO)); // idxBase @@ -139,20 +171,40 @@ inline cusparseStatus_t mat_vec_mul(cublasOperation_t transA, const sparse_mm &A */ cusparseStatus_t status; - status = cusparseScsrmv(A.cs_handle, - CUSPARSE_OPERATION_NON_TRANSPOSE, - m, - n, - A.A_csr.nnz, - &alpha, - A.descrA, - mA->val(), - mA->ptr(), - mA->ind(), - x, + // status = cusparseScsrmv(A.cs_handle, + // CUSPARSE_OPERATION_NON_TRANSPOSE, + // m, + // n, + // A.A_csr.nnz, + // &alpha, + // A.descrA, + // mA->val(), + // mA->ptr(), + // mA->ind(), + // x, +// &beta, + // y); + cusparseSpMatDescr_t matA; + cusparseCreateCsr(&matA,m,n,A.A_csr.nnz,(void*)mA->ptr(),(void*)mA->ind(),(void*)mA->val(),CUSPARSE_INDEX_32I,CUSPARSE_INDEX_32I,CUSPARSE_INDEX_BASE_ZERO,CUDA_R_32F); + cusparseDnVecDescr_t vecX,vecY; + cusparseCreateDnVec(&vecX,n, (void*)x, cudaDataType::CUDA_R_32F); + cusparseCreateDnVec(&vecY,m, (void*)y, cudaDataType::CUDA_R_32F); + unsigned long bufferSize = 0; + status = cusparseSpMV_bufferSize(A.cs_handle, CUSPARSE_OPERATION_NON_TRANSPOSE,&alpha,matA, + vecX, &beta, - y); - + vecY, + cudaDataType::CUDA_R_32F, + CUSPARSE_CSRMV_ALG1, &bufferSize); + void* buffer = NULL; + cudaMalloc(&buffer,bufferSize); + status = cusparseSpMV(A.cs_handle, CUSPARSE_OPERATION_NON_TRANSPOSE, + &alpha, matA, vecX, &beta, vecY, + CUDA_R_32F, CUSPARSE_CSRMV_ALG1, buffer); + cusparseDestroySpMat(matA); + cusparseDestroyDnVec(vecX); + cusparseDestroyDnVec(vecY); + cudaFree(buffer); return status; }