inline cusparseStatus_t cusparseTcsrmv (
cusparseHandle_t handle,
cusparseOperation_t transA,
int m,
int n,
int nnz,
const float *alpha,
const cusparseMatDescr_t descrA,
const float *csrValA,
const int *csrRowPtrA,
const int *csrColIndA,
const float *x,
const float *beta,
float *y
)
{
return cusparseScsrmv (
handle,
transA,
m,
n,
nnz,
alpha,
descrA,
csrValA,
csrRowPtrA,
csrColIndA,
x,
beta,
y
);
}
Multiply( const matrix_type & A ,
const size_type nrow ,
const size_type ncol ,
const vector_type & x ,
const vector_type & y )
{
CudaSparseSingleton & s = CudaSparseSingleton::singleton();
const scalar_type alpha = 1 , beta = 0 ;
cusparseStatus_t status =
cusparseScsrmv( s.handle ,
CUSPARSE_OPERATION_NON_TRANSPOSE ,
nrow , ncol , A.coefficients.dimension_0() ,
&alpha ,
s.descra ,
A.coefficients.ptr_on_device() ,
A.graph.row_map.ptr_on_device() ,
A.graph.entries.ptr_on_device() ,
x.ptr_on_device() ,
&beta ,
y.ptr_on_device() );
if ( CUSPARSE_STATUS_SUCCESS != status ) {
throw std::runtime_error( std::string("ERROR - cusparseDcsrmv " ) );
}
}
static void apply( const matrix_type & A ,
const vector_type & x ,
const vector_type & y )
{
CudaSparseSingleton & s = CudaSparseSingleton::singleton();
const float alpha = 1 , beta = 0 ;
const int n = A.graph.row_map.dimension_0() - 1 ;
const int nz = A.graph.entries.dimension_0();
cusparseStatus_t status =
cusparseScsrmv( s.handle ,
CUSPARSE_OPERATION_NON_TRANSPOSE ,
n , n , nz ,
&alpha ,
s.descra ,
A.values.ptr_on_device() ,
A.graph.row_map.ptr_on_device() ,
A.graph.entries.ptr_on_device() ,
x.ptr_on_device() ,
&beta ,
y.ptr_on_device() );
if ( CUSPARSE_STATUS_SUCCESS != status ) {
throw std::runtime_error( std::string("ERROR - cusparseScsrmv " ) );
}
}
void mul(const device_vector<float> &x, device_vector<float> &y,
float alpha = 1, bool append = false) const
{
float beta = append ? 1.0f : 0.0f;
cuda_check(
cusparseScsrmv(handle, CUSPARSE_OPERATION_NON_TRANSPOSE,
n, m, nnz, &alpha, desc.get(),
val.raw_ptr(), row.raw_ptr(), col.raw_ptr(),
x.raw_ptr(), &beta, y.raw_ptr()
)
);
}
int main(int argc, char **argv)
{
int N = 0, nz = 0, *I = NULL, *J = NULL;
float *val = NULL;
const float tol = 1e-5f;
const int max_iter = 10000;
float *x;
float *rhs;
float a, b, na, r0, r1;
float dot;
float *r, *p, *Ax;
int k;
float alpha, beta, alpham1;
printf("Starting [%s]...\n", sSDKname);
// This will pick the best possible CUDA capable device
cudaDeviceProp deviceProp;
int devID = findCudaDevice(argc, (const char **)argv);
checkCudaErrors(cudaGetDeviceProperties(&deviceProp, devID));
#if defined(__APPLE__) || defined(MACOSX)
fprintf(stderr, "Unified Memory not currently supported on OS X\n");
cudaDeviceReset();
exit(EXIT_WAIVED);
#endif
if (sizeof(void *) != 8)
{
fprintf(stderr, "Unified Memory requires compiling for a 64-bit system.\n");
cudaDeviceReset();
exit(EXIT_WAIVED);
}
if (((deviceProp.major << 4) + deviceProp.minor) < 0x30)
{
fprintf(stderr, "%s requires Compute Capability of SM 3.0 or higher to run.\nexiting...\n", argv[0]);
cudaDeviceReset();
exit(EXIT_WAIVED);
}
// Statistics about the GPU device
printf("> GPU device has %d Multi-Processors, SM %d.%d compute capabilities\n\n",
deviceProp.multiProcessorCount, deviceProp.major, deviceProp.minor);
/* Generate a random tridiagonal symmetric matrix in CSR format */
N = 1048576;
nz = (N-2)*3 + 4;
cudaMallocManaged((void **)&I, sizeof(int)*(N+1));
cudaMallocManaged((void **)&J, sizeof(int)*nz);
cudaMallocManaged((void **)&val, sizeof(float)*nz);
genTridiag(I, J, val, N, nz);
cudaMallocManaged((void **)&x, sizeof(float)*N);
cudaMallocManaged((void **)&rhs, sizeof(float)*N);
for (int i = 0; i < N; i++)
{
rhs[i] = 1.0;
x[i] = 0.0;
}
/* Get handle to the CUBLAS context */
cublasHandle_t cublasHandle = 0;
cublasStatus_t cublasStatus;
cublasStatus = cublasCreate(&cublasHandle);
checkCudaErrors(cublasStatus);
/* Get handle to the CUSPARSE context */
cusparseHandle_t cusparseHandle = 0;
cusparseStatus_t cusparseStatus;
cusparseStatus = cusparseCreate(&cusparseHandle);
checkCudaErrors(cusparseStatus);
cusparseMatDescr_t descr = 0;
cusparseStatus = cusparseCreateMatDescr(&descr);
checkCudaErrors(cusparseStatus);
cusparseSetMatType(descr,CUSPARSE_MATRIX_TYPE_GENERAL);
cusparseSetMatIndexBase(descr,CUSPARSE_INDEX_BASE_ZERO);
// temp memory for CG
checkCudaErrors(cudaMallocManaged((void **)&r, N*sizeof(float)));
checkCudaErrors(cudaMallocManaged((void **)&p, N*sizeof(float)));
checkCudaErrors(cudaMallocManaged((void **)&Ax, N*sizeof(float)));
cudaDeviceSynchronize();
for (int i=0; i < N; i++)
{
r[i] = rhs[i];
}
alpha = 1.0;
alpham1 = -1.0;
beta = 0.0;
r0 = 0.;
cusparseScsrmv(cusparseHandle,CUSPARSE_OPERATION_NON_TRANSPOSE, N, N, nz, &alpha, descr, val, I, J, x, &beta, Ax);
cublasSaxpy(cublasHandle, N, &alpham1, Ax, 1, r, 1);
cublasStatus = cublasSdot(cublasHandle, N, r, 1, r, 1, &r1);
k = 1;
while (r1 > tol*tol && k <= max_iter)
{
if (k > 1)
{
b = r1 / r0;
cublasStatus = cublasSscal(cublasHandle, N, &b, p, 1);
cublasStatus = cublasSaxpy(cublasHandle, N, &alpha, r, 1, p, 1);
}
else
{
cublasStatus = cublasScopy(cublasHandle, N, r, 1, p, 1);
}
cusparseScsrmv(cusparseHandle, CUSPARSE_OPERATION_NON_TRANSPOSE, N, N, nz, &alpha, descr, val, I, J, p, &beta, Ax);
cublasStatus = cublasSdot(cublasHandle, N, p, 1, Ax, 1, &dot);
a = r1 / dot;
cublasStatus = cublasSaxpy(cublasHandle, N, &a, p, 1, x, 1);
na = -a;
cublasStatus = cublasSaxpy(cublasHandle, N, &na, Ax, 1, r, 1);
r0 = r1;
cublasStatus = cublasSdot(cublasHandle, N, r, 1, r, 1, &r1);
cudaThreadSynchronize();
printf("iteration = %3d, residual = %e\n", k, sqrt(r1));
k++;
}
printf("Final residual: %e\n",sqrt(r1));
fprintf(stdout,"&&&& uvm_cg test %s\n", (sqrt(r1) < tol) ? "PASSED" : "FAILED");
float rsum, diff, err = 0.0;
for (int i = 0; i < N; i++)
{
rsum = 0.0;
for (int j = I[i]; j < I[i+1]; j++)
{
rsum += val[j]*x[J[j]];
}
diff = fabs(rsum - rhs[i]);
if (diff > err)
{
err = diff;
}
}
cusparseDestroy(cusparseHandle);
cublasDestroy(cublasHandle);
cudaFree(I);
cudaFree(J);
cudaFree(val);
cudaFree(x);
cudaFree(r);
cudaFree(p);
cudaFree(Ax);
cudaDeviceReset();
printf("Test Summary: Error amount = %f, result = %s\n", err, (k <= max_iter) ? "SUCCESS" : "FAILURE");
exit((k <= max_iter) ? EXIT_SUCCESS : EXIT_FAILURE);
}
/* Solve Ax=b using the conjugate gradient method a) without any preconditioning, b) using an Incomplete Cholesky preconditioner and c) using an ILU0 preconditioner. */
int main(int argc, char **argv)
{
const int max_iter = 1000;
int k, M = 0, N = 0, nz = 0, *I = NULL, *J = NULL;
int *d_col, *d_row;
int qatest = 0;
const float tol = 1e-12f;
float *x, *rhs;
float r0, r1, alpha, beta;
float *d_val, *d_x;
float *d_zm1, *d_zm2, *d_rm2;
float *d_r, *d_p, *d_omega, *d_y;
float *val = NULL;
float *d_valsILU0;
float *valsILU0;
float rsum, diff, err = 0.0;
float qaerr1, qaerr2 = 0.0;
float dot, numerator, denominator, nalpha;
const float floatone = 1.0;
const float floatzero = 0.0;
int nErrors = 0;
printf("conjugateGradientPrecond starting...\n");
/* QA testing mode */
if (checkCmdLineFlag(argc, (const char **)argv, "qatest"))
{
qatest = 1;
}
/* This will pick the best possible CUDA capable device */
cudaDeviceProp deviceProp;
int devID = findCudaDevice(argc, (const char **)argv);
printf("GPU selected Device ID = %d \n", devID);
if (devID < 0)
{
printf("Invalid GPU device %d selected, exiting...\n", devID);
exit(EXIT_SUCCESS);
}
checkCudaErrors(cudaGetDeviceProperties(&deviceProp, devID));
/* Statistics about the GPU device */
printf("> GPU device has %d Multi-Processors, SM %d.%d compute capabilities\n\n",
deviceProp.multiProcessorCount, deviceProp.major, deviceProp.minor);
int version = (deviceProp.major * 0x10 + deviceProp.minor);
if (version < 0x11)
{
printf("%s: requires a minimum CUDA compute 1.1 capability\n", sSDKname);
// cudaDeviceReset causes the driver to clean up all state. While
// not mandatory in normal operation, it is good practice. It is also
// needed to ensure correct operation when the application is being
// profiled. Calling cudaDeviceReset causes all profile data to be
// flushed before the application exits
cudaDeviceReset();
exit(EXIT_SUCCESS);
}
/* Generate a random tridiagonal symmetric matrix in CSR (Compressed Sparse Row) format */
M = N = 16384;
nz = 5*N-4*(int)sqrt((double)N);
I = (int *)malloc(sizeof(int)*(N+1)); // csr row pointers for matrix A
J = (int *)malloc(sizeof(int)*nz); // csr column indices for matrix A
val = (float *)malloc(sizeof(float)*nz); // csr values for matrix A
x = (float *)malloc(sizeof(float)*N);
rhs = (float *)malloc(sizeof(float)*N);
for (int i = 0; i < N; i++)
{
rhs[i] = 0.0; // Initialize RHS
x[i] = 0.0; // Initial approximation of solution
}
genLaplace(I, J, val, M, N, nz, rhs);
/* Create CUBLAS context */
cublasHandle_t cublasHandle = 0;
cublasStatus_t cublasStatus;
cublasStatus = cublasCreate(&cublasHandle);
checkCudaErrors(cublasStatus);
/* Create CUSPARSE context */
cusparseHandle_t cusparseHandle = 0;
cusparseStatus_t cusparseStatus;
cusparseStatus = cusparseCreate(&cusparseHandle);
checkCudaErrors(cusparseStatus);
/* Description of the A matrix*/
cusparseMatDescr_t descr = 0;
cusparseStatus = cusparseCreateMatDescr(&descr);
checkCudaErrors(cusparseStatus);
/* Define the properties of the matrix */
cusparseSetMatType(descr,CUSPARSE_MATRIX_TYPE_GENERAL);
cusparseSetMatIndexBase(descr,CUSPARSE_INDEX_BASE_ZERO);
/* Allocate required memory */
checkCudaErrors(cudaMalloc((void **)&d_col, nz*sizeof(int)));
checkCudaErrors(cudaMalloc((void **)&d_row, (N+1)*sizeof(int)));
checkCudaErrors(cudaMalloc((void **)&d_val, nz*sizeof(float)));
checkCudaErrors(cudaMalloc((void **)&d_x, N*sizeof(float)));
checkCudaErrors(cudaMalloc((void **)&d_y, N*sizeof(float)));
checkCudaErrors(cudaMalloc((void **)&d_r, N*sizeof(float)));
checkCudaErrors(cudaMalloc((void **)&d_p, N*sizeof(float)));
checkCudaErrors(cudaMalloc((void **)&d_omega, N*sizeof(float)));
cudaMemcpy(d_col, J, nz*sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(d_row, I, (N+1)*sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(d_val, val, nz*sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_r, rhs, N*sizeof(float), cudaMemcpyHostToDevice);
/* Conjugate gradient without preconditioning.
------------------------------------------
Follows the description by Golub & Van Loan, "Matrix Computations 3rd ed.", Section 10.2.6 */
printf("Convergence of conjugate gradient without preconditioning: \n");
k = 0;
r0 = 0;
cublasSdot(cublasHandle, N, d_r, 1, d_r, 1, &r1);
while (r1 > tol*tol && k <= max_iter)
{
k++;
if (k == 1)
{
cublasScopy(cublasHandle, N, d_r, 1, d_p, 1);
}
else
{
beta = r1/r0;
cublasSscal(cublasHandle, N, &beta, d_p, 1);
cublasSaxpy(cublasHandle, N, &floatone, d_r, 1, d_p, 1) ;
}
cusparseScsrmv(cusparseHandle,CUSPARSE_OPERATION_NON_TRANSPOSE, N, N, nz, &floatone, descr, d_val, d_row, d_col, d_p, &floatzero, d_omega);
cublasSdot(cublasHandle, N, d_p, 1, d_omega, 1, &dot);
alpha = r1/dot;
cublasSaxpy(cublasHandle, N, &alpha, d_p, 1, d_x, 1);
nalpha = -alpha;
cublasSaxpy(cublasHandle, N, &nalpha, d_omega, 1, d_r, 1);
r0 = r1;
cublasSdot(cublasHandle, N, d_r, 1, d_r, 1, &r1);
}
printf(" iteration = %3d, residual = %e \n", k, sqrt(r1));
cudaMemcpy(x, d_x, N*sizeof(float), cudaMemcpyDeviceToHost);
/* check result */
err = 0.0;
for (int i = 0; i < N; i++)
{
rsum = 0.0;
for (int j = I[i]; j < I[i+1]; j++)
{
rsum += val[j]*x[J[j]];
}
diff = fabs(rsum - rhs[i]);
if (diff > err)
{
err = diff;
}
}
printf(" Convergence Test: %s \n", (k <= max_iter) ? "OK" : "FAIL");
nErrors += (k > max_iter) ? 1 : 0;
qaerr1 = err;
if (0)
{
// output result in matlab-style array
int n=(int)sqrt((double)N);
printf("a = [ ");
for (int iy=0; iy<n; iy++)
{
for (int ix=0; ix<n; ix++)
{
printf(" %f ", x[iy*n+ix]);
}
if (iy == n-1)
{
printf(" ]");
}
printf("\n");
}
}
/* Preconditioned Conjugate Gradient using ILU.
--------------------------------------------
Follows the description by Golub & Van Loan, "Matrix Computations 3rd ed.", Algorithm 10.3.1 */
printf("\nConvergence of conjugate gradient using incomplete LU preconditioning: \n");
int nzILU0 = 2*N-1;
valsILU0 = (float *) malloc(nz*sizeof(float));
checkCudaErrors(cudaMalloc((void **)&d_valsILU0, nz*sizeof(float)));
checkCudaErrors(cudaMalloc((void **)&d_zm1, (N)*sizeof(float)));
checkCudaErrors(cudaMalloc((void **)&d_zm2, (N)*sizeof(float)));
checkCudaErrors(cudaMalloc((void **)&d_rm2, (N)*sizeof(float)));
/* create the analysis info object for the A matrix */
cusparseSolveAnalysisInfo_t infoA = 0;
cusparseStatus = cusparseCreateSolveAnalysisInfo(&infoA);
checkCudaErrors(cusparseStatus);
/* Perform the analysis for the Non-Transpose case */
cusparseStatus = cusparseScsrsv_analysis(cusparseHandle, CUSPARSE_OPERATION_NON_TRANSPOSE,
N, nz, descr, d_val, d_row, d_col, infoA);
checkCudaErrors(cusparseStatus);
/* Copy A data to ILU0 vals as input*/
cudaMemcpy(d_valsILU0, d_val, nz*sizeof(float), cudaMemcpyDeviceToDevice);
/* generate the Incomplete LU factor H for the matrix A using cudsparseScsrilu0 */
cusparseStatus = cusparseScsrilu0(cusparseHandle, CUSPARSE_OPERATION_NON_TRANSPOSE, N, descr, d_valsILU0, d_row, d_col, infoA);
checkCudaErrors(cusparseStatus);
/* Create info objects for the ILU0 preconditioner */
cusparseSolveAnalysisInfo_t info_u;
cusparseCreateSolveAnalysisInfo(&info_u);
cusparseMatDescr_t descrL = 0;
cusparseStatus = cusparseCreateMatDescr(&descrL);
cusparseSetMatType(descrL,CUSPARSE_MATRIX_TYPE_GENERAL);
cusparseSetMatIndexBase(descrL,CUSPARSE_INDEX_BASE_ZERO);
cusparseSetMatFillMode(descrL, CUSPARSE_FILL_MODE_LOWER);
cusparseSetMatDiagType(descrL, CUSPARSE_DIAG_TYPE_UNIT);
cusparseMatDescr_t descrU = 0;
cusparseStatus = cusparseCreateMatDescr(&descrU);
cusparseSetMatType(descrU,CUSPARSE_MATRIX_TYPE_GENERAL);
cusparseSetMatIndexBase(descrU,CUSPARSE_INDEX_BASE_ZERO);
cusparseSetMatFillMode(descrU, CUSPARSE_FILL_MODE_UPPER);
cusparseSetMatDiagType(descrU, CUSPARSE_DIAG_TYPE_NON_UNIT);
cusparseStatus = cusparseScsrsv_analysis(cusparseHandle, CUSPARSE_OPERATION_NON_TRANSPOSE, N, nz, descrU, d_val, d_row, d_col, info_u);
/* reset the initial guess of the solution to zero */
for (int i = 0; i < N; i++)
{
x[i] = 0.0;
}
checkCudaErrors(cudaMemcpy(d_r, rhs, N*sizeof(float), cudaMemcpyHostToDevice));
checkCudaErrors(cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyHostToDevice));
k = 0;
cublasSdot(cublasHandle, N, d_r, 1, d_r, 1, &r1);
while (r1 > tol*tol && k <= max_iter)
{
// Forward Solve, we can re-use infoA since the sparsity pattern of A matches that of L
cusparseStatus = cusparseScsrsv_solve(cusparseHandle, CUSPARSE_OPERATION_NON_TRANSPOSE, N, &floatone, descrL,
d_valsILU0, d_row, d_col, infoA, d_r, d_y);
checkCudaErrors(cusparseStatus);
// Back Substitution
cusparseStatus = cusparseScsrsv_solve(cusparseHandle, CUSPARSE_OPERATION_NON_TRANSPOSE, N, &floatone, descrU,
d_valsILU0, d_row, d_col, info_u, d_y, d_zm1);
checkCudaErrors(cusparseStatus);
k++;
if (k == 1)
{
cublasScopy(cublasHandle, N, d_zm1, 1, d_p, 1);
}
else
{
cublasSdot(cublasHandle, N, d_r, 1, d_zm1, 1, &numerator);
cublasSdot(cublasHandle, N, d_rm2, 1, d_zm2, 1, &denominator);
beta = numerator/denominator;
cublasSscal(cublasHandle, N, &beta, d_p, 1);
cublasSaxpy(cublasHandle, N, &floatone, d_zm1, 1, d_p, 1) ;
}
cusparseScsrmv(cusparseHandle,CUSPARSE_OPERATION_NON_TRANSPOSE, N, N, nzILU0, &floatone, descrU, d_val, d_row, d_col, d_p, &floatzero, d_omega);
cublasSdot(cublasHandle, N, d_r, 1, d_zm1, 1, &numerator);
cublasSdot(cublasHandle, N, d_p, 1, d_omega, 1, &denominator);
alpha = numerator / denominator;
cublasSaxpy(cublasHandle, N, &alpha, d_p, 1, d_x, 1);
cublasScopy(cublasHandle, N, d_r, 1, d_rm2, 1);
cublasScopy(cublasHandle, N, d_zm1, 1, d_zm2, 1);
nalpha = -alpha;
cublasSaxpy(cublasHandle, N, &nalpha, d_omega, 1, d_r, 1);
cublasSdot(cublasHandle, N, d_r, 1, d_r, 1, &r1);
}
printf(" iteration = %3d, residual = %e \n", k, sqrt(r1));
cudaMemcpy(x, d_x, N*sizeof(float), cudaMemcpyDeviceToHost);
/* check result */
err = 0.0;
for (int i = 0; i < N; i++)
{
rsum = 0.0;
for (int j = I[i]; j < I[i+1]; j++)
{
rsum += val[j]*x[J[j]];
}
diff = fabs(rsum - rhs[i]);
if (diff > err)
{
err = diff;
}
}
printf(" Convergence Test: %s \n", (k <= max_iter) ? "OK" : "FAIL");
nErrors += (k > max_iter) ? 1 : 0;
qaerr2 = err;
/* Destroy parameters */
cusparseDestroySolveAnalysisInfo(infoA);
cusparseDestroySolveAnalysisInfo(info_u);
/* Destroy contexts */
cusparseDestroy(cusparseHandle);
cublasDestroy(cublasHandle);
/* Free device memory */
free(I);
free(J);
free(val);
free(x);
free(rhs);
free(valsILU0);
cudaFree(d_col);
cudaFree(d_row);
cudaFree(d_val);
cudaFree(d_x);
cudaFree(d_y);
cudaFree(d_r);
cudaFree(d_p);
cudaFree(d_omega);
cudaFree(d_valsILU0);
cudaFree(d_zm1);
cudaFree(d_zm2);
cudaFree(d_rm2);
// cudaDeviceReset causes the driver to clean up all state. While
// not mandatory in normal operation, it is good practice. It is also
// needed to ensure correct operation when the application is being
// profiled. Calling cudaDeviceReset causes all profile data to be
// flushed before the application exits
cudaDeviceReset();
printf(" Test Summary:\n");
printf(" Counted total of %d errors\n", nErrors);
printf(" qaerr1 = %f qaerr2 = %f\n\n", fabs(qaerr1), fabs(qaerr2));
exit((nErrors == 0 &&fabs(qaerr1)<1e-5 && fabs(qaerr2) < 1e-5 ? EXIT_SUCCESS : EXIT_FAILURE));
}
void CQuadraticPath::cudaSolver(float* A, int* rowindex, int* columns,int N,int nz,float*Bx, float*X)
{
const int max_iter = 10000;
const float tol = 1e-12f;
float r0, r1, alpha, beta;
int *d_col, *d_row;
float *d_val, *d_x;
float *d_r, *d_p, *d_omega;
const float floatone = 1.0;
const float floatzero = 0.0;
float dot, nalpha;
/* Create CUBLAS context */
cublasHandle_t cublasHandle = 0;
cublasStatus_t cublasStatus;
cublasStatus = cublasCreate(&cublasHandle);
checkCudaErrors(cublasStatus);
/* Create CUSPARSE context */
cusparseHandle_t cusparseHandle = 0;
cusparseStatus_t cusparseStatus;
cusparseStatus = cusparseCreate(&cusparseHandle);
checkCudaErrors(cusparseStatus);
/* Description of the A matrix*/
cusparseMatDescr_t descr = 0;
cusparseStatus = cusparseCreateMatDescr(&descr);
checkCudaErrors(cusparseStatus);
/* Define the properties of the matrix */
cusparseSetMatType(descr,CUSPARSE_MATRIX_TYPE_GENERAL);
cusparseSetMatIndexBase(descr,CUSPARSE_INDEX_BASE_ZERO);
/* Allocate required memory */
checkCudaErrors(cudaMalloc((void **)&d_col, nz*sizeof(int)));
checkCudaErrors(cudaMalloc((void **)&d_row, (N+1)*sizeof(int)));
checkCudaErrors(cudaMalloc((void **)&d_val, nz*sizeof(float)));
checkCudaErrors(cudaMalloc((void **)&d_x, N*sizeof(float)));
checkCudaErrors(cudaMalloc((void **)&d_r, N*sizeof(float)));
checkCudaErrors(cudaMalloc((void **)&d_p, N*sizeof(float)));
checkCudaErrors(cudaMalloc((void **)&d_omega, N*sizeof(float)));
cudaMemcpy(d_col, columns, nz*sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(d_row, rowindex, (N+1)*sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(d_val, A, nz*sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_x, X, N*sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_r, Bx, N*sizeof(float), cudaMemcpyHostToDevice);
/* Conjugate gradient without preconditioning.
------------------------------------------
Follows the description by Golub & Van Loan, "Matrix Computations 3rd ed.", Section 10.2.6 */
int k = 0;
r0 = 0;
cublasSdot(cublasHandle, N, d_r, 1, d_r, 1, &r1);
while (r1 > tol*tol && k <= max_iter)
{
k++;
if (k == 1)
{
cublasScopy(cublasHandle, N, d_r, 1, d_p, 1);
}
else
{
beta = r1/r0;
cublasSscal(cublasHandle, N, &beta, d_p, 1);
cublasSaxpy(cublasHandle, N, &floatone, d_r, 1, d_p, 1) ;
}
cusparseScsrmv(cusparseHandle,CUSPARSE_OPERATION_NON_TRANSPOSE, N, N, nz, &floatone, descr, d_val, d_row, d_col, d_p, &floatzero, d_omega);
cublasSdot(cublasHandle, N, d_p, 1, d_omega, 1, &dot);
alpha = r1/dot;
cublasSaxpy(cublasHandle, N, &alpha, d_p, 1, d_x, 1);
nalpha = -alpha;
cublasSaxpy(cublasHandle, N, &nalpha, d_omega, 1, d_r, 1);
r0 = r1;
cublasSdot(cublasHandle, N, d_r, 1, d_r, 1, &r1);
}
cudaMemcpy(X, d_x, N*sizeof(float), cudaMemcpyDeviceToHost);
cudaFree(d_col);
cudaFree(d_row);
cudaFree(d_val);
cudaFree(d_x);
cudaFree(d_r);
cudaFree(d_p);
cudaFree(d_omega);
}
extern "C" magma_int_t
magma_s_spmv(
float alpha,
magma_s_sparse_matrix A,
magma_s_vector x,
float beta,
magma_s_vector y,
magma_queue_t queue )
{
// set queue for old dense routines
magma_queue_t orig_queue;
magmablasGetKernelStream( &orig_queue );
if ( A.memory_location != x.memory_location ||
x.memory_location != y.memory_location ) {
printf("error: linear algebra objects are not located in same memory!\n");
printf("memory locations are: %d %d %d\n",
A.memory_location, x.memory_location, y.memory_location );
magmablasSetKernelStream( orig_queue );
return MAGMA_ERR_INVALID_PTR;
}
// DEV case
if ( A.memory_location == Magma_DEV ) {
if ( A.num_cols == x.num_rows && x.num_cols == 1 ) {
if ( A.storage_type == Magma_CSR
|| A.storage_type == Magma_CSRL
|| A.storage_type == Magma_CSRU ) {
//printf("using CSR kernel for SpMV: ");
//magma_sgecsrmv( MagmaNoTrans, A.num_rows, A.num_cols, alpha,
// A.dval, A.drow, A.dcol, x.dval, beta, y.dval );
//printf("done.\n");
cusparseHandle_t cusparseHandle = 0;
cusparseStatus_t cusparseStatus;
cusparseStatus = cusparseCreate(&cusparseHandle);
cusparseSetStream( cusparseHandle, queue );
cusparseMatDescr_t descr = 0;
cusparseStatus = cusparseCreateMatDescr(&descr);
cusparseSetMatType(descr,CUSPARSE_MATRIX_TYPE_GENERAL);
cusparseSetMatIndexBase(descr,CUSPARSE_INDEX_BASE_ZERO);
cusparseScsrmv( cusparseHandle,CUSPARSE_OPERATION_NON_TRANSPOSE,
A.num_rows, A.num_cols, A.nnz, &alpha, descr,
A.dval, A.drow, A.dcol, x.dval, &beta, y.dval );
cusparseDestroyMatDescr( descr );
cusparseDestroy( cusparseHandle );
}
else if ( A.storage_type == Magma_ELL ) {
//printf("using ELLPACKT kernel for SpMV: ");
magma_sgeelltmv( MagmaNoTrans, A.num_rows, A.num_cols,
A.max_nnz_row, alpha, A.dval, A.dcol, x.dval, beta,
y.dval, queue );
//printf("done.\n");
}
else if ( A.storage_type == Magma_ELLPACKT ) {
//printf("using ELL kernel for SpMV: ");
magma_sgeellmv( MagmaNoTrans, A.num_rows, A.num_cols,
A.max_nnz_row, alpha, A.dval, A.dcol, x.dval, beta,
y.dval, queue );
//printf("done.\n");
}
else if ( A.storage_type == Magma_ELLRT ) {
//printf("using ELLRT kernel for SpMV: ");
magma_sgeellrtmv( MagmaNoTrans, A.num_rows, A.num_cols,
A.max_nnz_row, alpha, A.dval, A.dcol, A.drow, x.dval,
beta, y.dval, A.alignment, A.blocksize, queue );
//printf("done.\n");
}
else if ( A.storage_type == Magma_SELLP ) {
//printf("using SELLP kernel for SpMV: ");
magma_sgesellpmv( MagmaNoTrans, A.num_rows, A.num_cols,
A.blocksize, A.numblocks, A.alignment,
alpha, A.dval, A.dcol, A.drow, x.dval, beta, y.dval, queue );
//printf("done.\n");
}
else if ( A.storage_type == Magma_DENSE ) {
//printf("using DENSE kernel for SpMV: ");
magmablas_sgemv( MagmaNoTrans, A.num_rows, A.num_cols, alpha,
A.dval, A.num_rows, x.dval, 1, beta, y.dval,
1 );
//printf("done.\n");
}
/* else if ( A.storage_type == Magma_BCSR ) {
//printf("using CUSPARSE BCSR kernel for SpMV: ");
// CUSPARSE context //
cusparseHandle_t cusparseHandle = 0;
cusparseStatus_t cusparseStatus;
cusparseStatus = cusparseCreate(&cusparseHandle);
cusparseSetStream( cusparseHandle, queue );
cusparseMatDescr_t descr = 0;
cusparseStatus = cusparseCreateMatDescr(&descr);
// end CUSPARSE context //
cusparseDirection_t dirA = CUSPARSE_DIRECTION_ROW;
int mb = (A.num_rows + A.blocksize-1)/A.blocksize;
int nb = (A.num_cols + A.blocksize-1)/A.blocksize;
cusparseSbsrmv( cusparseHandle, dirA,
CUSPARSE_OPERATION_NON_TRANSPOSE, mb, nb, A.numblocks,
&alpha, descr, A.dval, A.drow, A.dcol, A.blocksize, x.dval,
&beta, y.dval );
//printf("done.\n");
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
}*/
else {
printf("error: format not supported.\n");
magmablasSetKernelStream( orig_queue );
return MAGMA_ERR_NOT_SUPPORTED;
}
}
else if ( A.num_cols < x.num_rows || x.num_cols > 1 ) {
magma_int_t num_vecs = x.num_rows / A.num_cols * x.num_cols;
if ( A.storage_type == Magma_CSR ) {
cusparseHandle_t cusparseHandle = 0;
cusparseStatus_t cusparseStatus;
cusparseStatus = cusparseCreate(&cusparseHandle);
cusparseSetStream( cusparseHandle, queue );
cusparseMatDescr_t descr = 0;
cusparseStatus = cusparseCreateMatDescr(&descr);
cusparseSetMatType(descr,CUSPARSE_MATRIX_TYPE_GENERAL);
cusparseSetMatIndexBase(descr,CUSPARSE_INDEX_BASE_ZERO);
if ( x.major == MagmaColMajor) {
cusparseScsrmm(cusparseHandle,
CUSPARSE_OPERATION_NON_TRANSPOSE,
A.num_rows, num_vecs, A.num_cols, A.nnz,
&alpha, descr, A.dval, A.drow, A.dcol,
x.dval, A.num_cols, &beta, y.dval, A.num_cols);
} else if ( x.major == MagmaRowMajor) {
cusparseScsrmm2(cusparseHandle,
CUSPARSE_OPERATION_NON_TRANSPOSE,
CUSPARSE_OPERATION_TRANSPOSE,
A.num_rows, num_vecs, A.num_cols, A.nnz,
&alpha, descr, A.dval, A.drow, A.dcol,
x.dval, A.num_cols, &beta, y.dval, A.num_cols);
}
cusparseDestroyMatDescr( descr );
cusparseDestroy( cusparseHandle );
}
else if ( A.storage_type == Magma_ELL ) {
if ( x.major == MagmaColMajor) {
magma_smgeelltmv( MagmaNoTrans, A.num_rows, A.num_cols,
num_vecs, A.max_nnz_row, alpha, A.dval, A.dcol, x.dval,
beta, y.dval, queue );
}
else if ( x.major == MagmaRowMajor) {
// transpose first to col major
magma_s_vector x2;
magma_svtranspose( x, &x2, queue );
magma_smgeellmv( MagmaNoTrans, A.num_rows, A.num_cols,
num_vecs, A.max_nnz_row, alpha, A.dval, A.dcol, x.dval,
beta, y.dval, queue );
magma_s_vfree(&x2, queue );
}
}
else if ( A.storage_type == Magma_ELLPACKT ) {
if ( x.major == MagmaColMajor) {
magma_smgeellmv( MagmaNoTrans, A.num_rows, A.num_cols,
num_vecs, A.max_nnz_row, alpha, A.dval, A.dcol, x.dval,
beta, y.dval, queue );
}
else if ( x.major == MagmaRowMajor) {
// transpose first to col major
magma_s_vector x2;
magma_svtranspose( x, &x2, queue );
magma_smgeelltmv( MagmaNoTrans, A.num_rows, A.num_cols,
num_vecs, A.max_nnz_row, alpha, A.dval, A.dcol, x.dval,
beta, y.dval, queue );
magma_s_vfree(&x2, queue );
}
} else if ( A.storage_type == Magma_SELLP ) {
if ( x.major == MagmaRowMajor) {
magma_smgesellpmv( MagmaNoTrans, A.num_rows, A.num_cols,
num_vecs, A.blocksize, A.numblocks, A.alignment,
alpha, A.dval, A.dcol, A.drow, x.dval, beta, y.dval, queue );
}
else if ( x.major == MagmaColMajor) {
// transpose first to row major
magma_s_vector x2;
magma_svtranspose( x, &x2, queue );
magma_smgesellpmv( MagmaNoTrans, A.num_rows, A.num_cols,
num_vecs, A.blocksize, A.numblocks, A.alignment,
alpha, A.dval, A.dcol, A.drow, x2.dval, beta, y.dval, queue );
magma_s_vfree(&x2, queue );
}
}/*
if ( A.storage_type == Magma_DENSE ) {
//printf("using DENSE kernel for SpMV: ");
magmablas_smgemv( MagmaNoTrans, A.num_rows, A.num_cols,
num_vecs, alpha, A.dval, A.num_rows, x.dval, 1,
beta, y.dval, 1 );
//printf("done.\n");
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
}*/
else {
printf("error: format not supported.\n");
magmablasSetKernelStream( orig_queue );
return MAGMA_ERR_NOT_SUPPORTED;
}
}
}
// CPU case missing!
else {
printf("error: CPU not yet supported.\n");
magmablasSetKernelStream( orig_queue );
return MAGMA_ERR_NOT_SUPPORTED;
}
magmablasSetKernelStream( orig_queue );
return MAGMA_SUCCESS;
}
int main(int argc, char **argv)
{
int M = 0, N = 0, nz = 0, *I = NULL, *J = NULL;
float *val = NULL;
const float tol = 1e-5f;
const int max_iter = 10000;
float *x;
float *rhs;
float a, b, na, r0, r1;
int *d_col, *d_row;
float *d_val, *d_x, dot;
float *d_r, *d_p, *d_Ax;
int k;
float alpha, beta, alpham1;
shrQAStart(argc, argv);
// This will pick the best possible CUDA capable device
cudaDeviceProp deviceProp;
int devID = findCudaDevice(argc, (const char **)argv);
if (devID < 0) {
printf("exiting...\n");
shrQAFinishExit(argc, (const char **)argv, QA_PASSED);
exit(0);
}
checkCudaErrors( cudaGetDeviceProperties(&deviceProp, devID) );
// Statistics about the GPU device
printf("> GPU device has %d Multi-Processors, SM %d.%d compute capabilities\n\n",
deviceProp.multiProcessorCount, deviceProp.major, deviceProp.minor);
int version = (deviceProp.major * 0x10 + deviceProp.minor);
if(version < 0x11)
{
printf("%s: requires a minimum CUDA compute 1.1 capability\n", sSDKname);
cudaDeviceReset();
shrQAFinishExit(argc, (const char **)argv, QA_PASSED);
}
/* Generate a random tridiagonal symmetric matrix in CSR format */
M = N = 1048576;
nz = (N-2)*3 + 4;
I = (int*)malloc(sizeof(int)*(N+1));
J = (int*)malloc(sizeof(int)*nz);
val = (float*)malloc(sizeof(float)*nz);
genTridiag(I, J, val, N, nz);
x = (float*)malloc(sizeof(float)*N);
rhs = (float*)malloc(sizeof(float)*N);
for (int i = 0; i < N; i++) {
rhs[i] = 1.0;
x[i] = 0.0;
}
/* Get handle to the CUBLAS context */
cublasHandle_t cublasHandle = 0;
cublasStatus_t cublasStatus;
cublasStatus = cublasCreate(&cublasHandle);
if ( checkCublasStatus (cublasStatus, "!!!! CUBLAS initialization error\n") ) return EXIT_FAILURE;
/* Get handle to the CUSPARSE context */
cusparseHandle_t cusparseHandle = 0;
cusparseStatus_t cusparseStatus;
cusparseStatus = cusparseCreate(&cusparseHandle);
if ( checkCusparseStatus (cusparseStatus, "!!!! CUSPARSE initialization error\n") ) return EXIT_FAILURE;
cusparseMatDescr_t descr = 0;
cusparseStatus = cusparseCreateMatDescr(&descr);
if ( checkCusparseStatus (cusparseStatus, "!!!! CUSPARSE cusparseCreateMatDescr error\n") ) return EXIT_FAILURE;
cusparseSetMatType(descr,CUSPARSE_MATRIX_TYPE_GENERAL);
cusparseSetMatIndexBase(descr,CUSPARSE_INDEX_BASE_ZERO);
checkCudaErrors( cudaMalloc((void**)&d_col, nz*sizeof(int)) );
checkCudaErrors( cudaMalloc((void**)&d_row, (N+1)*sizeof(int)) );
checkCudaErrors( cudaMalloc((void**)&d_val, nz*sizeof(float)) );
checkCudaErrors( cudaMalloc((void**)&d_x, N*sizeof(float)) );
checkCudaErrors( cudaMalloc((void**)&d_r, N*sizeof(float)) );
checkCudaErrors( cudaMalloc((void**)&d_p, N*sizeof(float)) );
checkCudaErrors( cudaMalloc((void**)&d_Ax, N*sizeof(float)) );
cudaMemcpy(d_col, J, nz*sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(d_row, I, (N+1)*sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(d_val, val, nz*sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_r, rhs, N*sizeof(float), cudaMemcpyHostToDevice);
alpha = 1.0;
alpham1 = -1.0;
beta = 0.0;
r0 = 0.;
cusparseScsrmv(cusparseHandle,CUSPARSE_OPERATION_NON_TRANSPOSE, N, N, nz, &alpha, descr, d_val, d_row, d_col, d_x, &beta, d_Ax);
cublasSaxpy(cublasHandle, N, &alpham1, d_Ax, 1, d_r, 1);
cublasStatus = cublasSdot(cublasHandle, N, d_r, 1, d_r, 1, &r1);
k = 1;
while (r1 > tol*tol && k <= max_iter) {
if (k > 1) {
b = r1 / r0;
cublasStatus = cublasSscal(cublasHandle, N, &b, d_p, 1);
cublasStatus = cublasSaxpy(cublasHandle, N, &alpha, d_r, 1, d_p, 1);
} else {
cublasStatus = cublasScopy(cublasHandle, N, d_r, 1, d_p, 1);
}
cusparseScsrmv(cusparseHandle, CUSPARSE_OPERATION_NON_TRANSPOSE, N, N, nz, &alpha, descr, d_val, d_row, d_col, d_p, &beta, d_Ax);
cublasStatus = cublasSdot(cublasHandle, N, d_p, 1, d_Ax, 1, &dot);
a = r1 / dot;
cublasStatus = cublasSaxpy(cublasHandle, N, &a, d_p, 1, d_x, 1);
na = -a;
cublasStatus = cublasSaxpy(cublasHandle, N, &na, d_Ax, 1, d_r, 1);
r0 = r1;
cublasStatus = cublasSdot(cublasHandle, N, d_r, 1, d_r, 1, &r1);
cudaThreadSynchronize();
printf("iteration = %3d, residual = %e\n", k, sqrt(r1));
k++;
}
cudaMemcpy(x, d_x, N*sizeof(float), cudaMemcpyDeviceToHost);
float rsum, diff, err = 0.0;
for (int i = 0; i < N; i++) {
rsum = 0.0;
for (int j = I[i]; j < I[i+1]; j++) {
rsum += val[j]*x[J[j]];
}
diff = fabs(rsum - rhs[i]);
if (diff > err) err = diff;
}
cusparseDestroy(cusparseHandle);
cublasDestroy(cublasHandle);
free(I);
free(J);
free(val);
free(x);
free(rhs);
cudaFree(d_col);
cudaFree(d_row);
cudaFree(d_val);
cudaFree(d_x);
cudaFree(d_r);
cudaFree(d_p);
cudaFree(d_Ax);
cudaDeviceReset();
printf("Test Summary: Error amount = %f\n", err);
shrQAFinishExit(argc, (const char **)argv, (k <= max_iter) ? QA_PASSED : QA_FAILED );
}
