fdurastante
/
psblas3
mirror of https://github.com/sfilippone/psblas3.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

Fix usage of HAVE_CUDA/HAVE_GPU (mostly disappeared)

Browse Source
repack-nvid
sfilippone 1 year ago
parent e373ed7e0b
commit b2b7b074df
376 changed files with 26 additions and 3547 deletions
Show Stats Download Patch File Download Diff File

4
cuda/base_cusparse_mod.F90
Unescape Escape View File

@ -79,9 +79,6 @@ module base_cusparse_mod
enumerator cusparse_direction_column
end enum
#if defined(HAVE_CUDA) && defined(HAVE_SPGPU)
interface
function FcusparseCreate() &
& bind(c,name="FcusparseCreate") result(res)
@ -113,5 +110,4 @@ contains
res = FcusparseDestroy()
end function closeFcusparse
#endif
end module base_cusparse_mod

5
cuda/c_cusparse_mod.F90
Unescape Escape View File

@ -43,9 +43,6 @@ module c_cusparse_mod
end type c_Hmat
#endif
#if defined(HAVE_CUDA) && defined(HAVE_SPGPU)
interface CSRGDeviceFree
function c_CSRGDeviceFree(Mat) &
& bind(c,name="c_CSRGDeviceFree") result(res)
@ -300,6 +297,4 @@ module c_cusparse_mod
end interface
#endif
#endif
end module c_cusparse_mod

2
cuda/ccusparse.c
Unescape Escape View File

@ -33,7 +33,6 @@
#include <stdio.h>
#include <stdlib.h>
#ifdef HAVE_SPGPU
#include <cuda_runtime.h>
#include <cusparse_v2.h>
#include "cintrf.h"
@ -94,4 +93,3 @@
#include "fcusparse_fct.h"
#endif

4
cuda/cintrf.h
Unescape Escape View File

@ -36,7 +36,6 @@
#include <stdlib.h>
#include <stdio.h>
#if defined(HAVE_SPGPU) && defined(HAVE_CUDA)
#include "core.h"
#include "cuda_util.h"
#include "vector.h"
@ -45,7 +44,4 @@
#define ELL_PITCH_ALIGN_S 32
#define ELL_PITCH_ALIGN_D 16
#endif
#endif

9
cuda/cuda_util.c
Unescape Escape View File

@ -32,8 +32,6 @@
#include "cuda_util.h"
#if defined(HAVE_CUDA)
static int hasUVA=-1;
static struct cudaDeviceProp *prop=NULL;
@ -440,10 +438,6 @@ void psb_cudaDestroyCublasHandle()
psb_cublas_handle=NULL;
}
/* Simple memory tools */
int allocateInt(void **d_int, int n)
@ -803,6 +797,3 @@ double etime()
}
#endif

2
cuda/cuda_util.h
Unescape Escape View File

@ -38,7 +38,6 @@
#include <sys/time.h>
#include <string.h>
#if defined(HAVE_CUDA)
#include "cuda_runtime.h"
#include "core.h"
#include "cuComplex.h"
@ -134,6 +133,5 @@ void freeDoubleComplex(void *);
double etime();
#endif
#endif

2
cuda/cvectordev.c
Unescape Escape View File

@ -32,7 +32,6 @@
#include <stdio.h>
#include <complex.h>
#if defined(HAVE_SPGPU)
//#include "utils.h"
//#include "common.h"
#include "cvectordev.h"
@ -321,5 +320,4 @@ int absMultiVecDeviceFloatComplex(int n, cuFloatComplex alpha, void *deviceVecA)
return(i);
}
#endif

3
cuda/cvectordev.h
Unescape Escape View File

@ -31,7 +31,6 @@
#pragma once
#if defined(HAVE_SPGPU)
//#include "utils.h"
#include <complex.h>
#include "cuComplex.h"
@ -77,5 +76,3 @@ int absMultiVecDeviceFloatComplex(int n, cuFloatComplex alpha, void *deviceVecA)
int absMultiVecDeviceFloatComplex2(int n, cuFloatComplex alpha,
void *deviceVecA, void *deviceVecB);
#endif

4
cuda/d_cusparse_mod.F90
Unescape Escape View File

@ -43,9 +43,6 @@ module d_cusparse_mod
end type d_Hmat
#endif
#if defined(HAVE_CUDA) && defined(HAVE_SPGPU)
interface CSRGDeviceFree
function d_CSRGDeviceFree(Mat) &
& bind(c,name="d_CSRGDeviceFree") result(res)
@ -298,7 +295,6 @@ module d_cusparse_mod
integer(c_int) :: res
end function d_HYBGHost2Device
end interface
#endif
#endif

2
cuda/dcusparse.c
Unescape Escape View File

@ -33,7 +33,6 @@
#include <stdio.h>
#include <stdlib.h>
#ifdef HAVE_SPGPU
#include <cuda_runtime.h>
#include <cusparse_v2.h>
#include "cintrf.h"
@ -92,4 +91,3 @@
#include "fcusparse_fct.h"
#endif

30
cuda/diagdev.c
Unescape Escape View File

@ -34,7 +34,6 @@
#include <sys/stat.h>
#include <fcntl.h>
#include <unistd.h>
#if defined(HAVE_SPGPU)
//new
DiagDeviceParams getDiagDeviceParams(unsigned int rows, unsigned int columns, unsigned int diags, unsigned int elementType)
{
@ -111,7 +110,6 @@ void freeDiagDevice(void* remoteMatrix)
//new
int FallocDiagDevice(void** deviceMat, unsigned int rows, unsigned int columns,unsigned int diags,unsigned int elementType)
{ int i;
#ifdef HAVE_SPGPU
DiagDeviceParams p;
p = getDiagDeviceParams(rows, columns, diags,elementType);
@ -120,15 +118,11 @@ int FallocDiagDevice(void** deviceMat, unsigned int rows, unsigned int columns,u
fprintf(stderr,"From routine : %s : %d \n","FallocEllDevice",i);
}
return(i);
#else
return SPGPU_UNSUPPORTED;
#endif
}
int writeDiagDeviceDouble(void* deviceMat, double* a, int* off, int n)
{ int i,fo,fa;
char buf_a[255], buf_o[255],tmp[255];
#ifdef HAVE_SPGPU
struct DiagDevice *devMat = (struct DiagDevice *) deviceMat;
// Ex updateFromHost function
/* memset(buf_a,'\0',255); */
@ -159,14 +153,10 @@ int writeDiagDeviceDouble(void* deviceMat, double* a, int* off, int n)
return SPGPU_SUCCESS;
else
return SPGPU_UNSUPPORTED;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int readDiagDeviceDouble(void* deviceMat, double* a, int* off)
{ int i;
#ifdef HAVE_SPGPU
struct DiagDevice *devMat = (struct DiagDevice *) deviceMat;
i = readRemoteBuffer((void *) a, (void *)devMat->cM,devMat->rows*devMat->diags*sizeof(double));
i = readRemoteBuffer((void *) off, (void *)devMat->off, devMat->diags*sizeof(int));
@ -174,9 +164,6 @@ int readDiagDeviceDouble(void* deviceMat, double* a, int* off)
fprintf(stderr,"From routine : %s : %d \n","readEllDeviceDouble",i);
}*/
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
//new
@ -188,7 +175,6 @@ int spmvDiagDeviceDouble(void *deviceMat, double alpha, void* deviceX,
struct MultiVectDevice *y = (struct MultiVectDevice *) deviceY;
spgpuHandle_t handle=psb_cudaGetHandle();
#ifdef HAVE_SPGPU
#ifdef VERBOSE
/*__assert(x->count_ == x->count_, "ERROR: x and y don't share the same number of vectors");*/
/*__assert(x->size_ >= devMat->columns, "ERROR: x vector's size is not >= to matrix size (columns)");*/
@ -201,16 +187,12 @@ int spmvDiagDeviceDouble(void *deviceMat, double alpha, void* deviceX,
//cudaSync();
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int writeDiagDeviceFloat(void* deviceMat, float* a, int* off, int n)
{ int i,fo,fa;
char buf_a[255], buf_o[255],tmp[255];
#ifdef HAVE_SPGPU
struct DiagDevice *devMat = (struct DiagDevice *) deviceMat;
// Ex updateFromHost function
/* memset(buf_a,'\0',255); */
@ -241,14 +223,10 @@ int writeDiagDeviceFloat(void* deviceMat, float* a, int* off, int n)
return SPGPU_SUCCESS;
else
return SPGPU_UNSUPPORTED;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int readDiagDeviceFloat(void* deviceMat, float* a, int* off)
{ int i;
#ifdef HAVE_SPGPU
struct DiagDevice *devMat = (struct DiagDevice *) deviceMat;
i = readRemoteBuffer((void *) a, (void *)devMat->cM,devMat->rows*devMat->diags*sizeof(float));
i = readRemoteBuffer((void *) off, (void *)devMat->off, devMat->diags*sizeof(int));
@ -256,9 +234,6 @@ int readDiagDeviceFloat(void* deviceMat, float* a, int* off)
fprintf(stderr,"From routine : %s : %d \n","readEllDeviceFloat",i);
}*/
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
//new
@ -270,7 +245,6 @@ int spmvDiagDeviceFloat(void *deviceMat, float alpha, void* deviceX,
struct MultiVectDevice *y = (struct MultiVectDevice *) deviceY;
spgpuHandle_t handle=psb_cudaGetHandle();
#ifdef HAVE_SPGPU
#ifdef VERBOSE
/*__assert(x->count_ == x->count_, "ERROR: x and y don't share the same number of vectors");*/
/*__assert(x->size_ >= devMat->columns, "ERROR: x vector's size is not >= to matrix size (columns)");*/
@ -283,9 +257,5 @@ int spmvDiagDeviceFloat(void *deviceMat, float alpha, void* deviceX,
//cudaSync();
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
#endif

5
cuda/diagdev.h
Unescape Escape View File

@ -32,7 +32,6 @@
#ifndef _DIAGDEV_H_
#define _DIAGDEV_H_
#ifdef HAVE_SPGPU
#include "cintrf.h"
#include "dia.h"
@ -88,8 +87,4 @@ int spmvDiagDeviceFloat(void *deviceMat, float alpha, void* deviceX,
#else
#define CINTRF_UNSUPPORTED -1
#endif
#endif

7
cuda/diagdev_mod.F90
Unescape Escape View File

@ -41,8 +41,6 @@ module diagdev_mod
integer(c_int) :: firstIndex
end type diagdev_parms
#ifdef HAVE_SPGPU
interface
function FgetDiagDeviceParams(rows, columns, elementType, firstIndex) &
& result(res) bind(c,name='getDiagDeviceParams')
@ -65,7 +63,6 @@ module diagdev_mod
end function FallocDiagDevice
end interface
interface writeDiagDevice
function writeDiagDeviceFloat(deviceMat,a,off,n) &
@ -174,7 +171,6 @@ module diagdev_mod
end function getDiagTimer
end interface
interface
function getDiagDevicePitch(deviceMat) &
& bind(c,name='getDiagDevicePitch') result(res)
@ -225,7 +221,4 @@ module diagdev_mod
end function spmvDiagDeviceDoubleComplex
end interface spmvDiagDevice
#endif
end module diagdev_mod

62
cuda/dnsdev.c
Unescape Escape View File

@ -31,8 +31,6 @@
#include <sys/time.h>
#include "dnsdev.h"
#if defined(HAVE_SPGPU)
#define PASS_RS 0
#define IMIN(a,b) ((a)<(b) ? (a) : (b))
@ -102,7 +100,6 @@ int FallocDnsDevice(void** deviceMat, unsigned int rows,
unsigned int columns, unsigned int elementType,
unsigned int firstIndex)
{ int i;
#ifdef HAVE_SPGPU
DnsDeviceParams p;
p = getDnsDeviceParams(rows, columns, elementType, firstIndex);
@ -111,9 +108,6 @@ int FallocDnsDevice(void** deviceMat, unsigned int rows,
fprintf(stderr,"From routine : %s : %d \n","FallocDnsDevice",i);
}
return(i);
#else
return SPGPU_UNSUPPORTED;
#endif
}
@ -124,7 +118,6 @@ int spmvDnsDeviceFloat(char transa, int m, int n, int k, float *alpha,
struct MultiVectDevice *x = (struct MultiVectDevice *) deviceX;
struct MultiVectDevice *y = (struct MultiVectDevice *) deviceY;
int status;
#ifdef HAVE_SPGPU
cublasHandle_t handle=psb_cudaGetCublasHandle();
cublasOperation_t trans=((transa == 'N')? CUBLAS_OP_N:((transa=='T')? CUBLAS_OP_T:CUBLAS_OP_C));
@ -143,9 +136,6 @@ int spmvDnsDeviceFloat(char transa, int m, int n, int k, float *alpha,
return SPGPU_SUCCESS;
else
return SPGPU_UNSUPPORTED;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int spmvDnsDeviceDouble(char transa, int m, int n, int k, double *alpha,
@ -155,7 +145,6 @@ int spmvDnsDeviceDouble(char transa, int m, int n, int k, double *alpha,
struct MultiVectDevice *x = (struct MultiVectDevice *) deviceX;
struct MultiVectDevice *y = (struct MultiVectDevice *) deviceY;
int status;
#ifdef HAVE_SPGPU
cublasHandle_t handle=psb_cudaGetCublasHandle();
cublasOperation_t trans=((transa == 'N')? CUBLAS_OP_N:((transa=='T')? CUBLAS_OP_T:CUBLAS_OP_C));
@ -174,9 +163,6 @@ int spmvDnsDeviceDouble(char transa, int m, int n, int k, double *alpha,
return SPGPU_SUCCESS;
else
return SPGPU_UNSUPPORTED;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int spmvDnsDeviceFloatComplex(char transa, int m, int n, int k, float complex *alpha,
@ -186,7 +172,6 @@ int spmvDnsDeviceFloatComplex(char transa, int m, int n, int k, float complex *a
struct MultiVectDevice *x = (struct MultiVectDevice *) deviceX;
struct MultiVectDevice *y = (struct MultiVectDevice *) deviceY;
int status;
#ifdef HAVE_SPGPU
cublasHandle_t handle=psb_cudaGetCublasHandle();
cublasOperation_t trans=((transa == 'N')? CUBLAS_OP_N:((transa=='T')? CUBLAS_OP_T:CUBLAS_OP_C));
@ -205,9 +190,6 @@ int spmvDnsDeviceFloatComplex(char transa, int m, int n, int k, float complex *a
return SPGPU_SUCCESS;
else
return SPGPU_UNSUPPORTED;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int spmvDnsDeviceDoubleComplex(char transa, int m, int n, int k, double complex *alpha,
@ -217,7 +199,6 @@ int spmvDnsDeviceDoubleComplex(char transa, int m, int n, int k, double complex
struct MultiVectDevice *x = (struct MultiVectDevice *) deviceX;
struct MultiVectDevice *y = (struct MultiVectDevice *) deviceY;
int status;
#ifdef HAVE_SPGPU
cublasHandle_t handle=psb_cudaGetCublasHandle();
cublasOperation_t trans=((transa == 'N')? CUBLAS_OP_N:((transa=='T')? CUBLAS_OP_T:CUBLAS_OP_C));
@ -236,15 +217,11 @@ int spmvDnsDeviceDoubleComplex(char transa, int m, int n, int k, double complex
return SPGPU_SUCCESS;
else
return SPGPU_UNSUPPORTED;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int writeDnsDeviceFloat(void* deviceMat, float* val, int lda, int nc)
{ int i;
#ifdef HAVE_SPGPU
struct DnsDevice *devMat = (struct DnsDevice *) deviceMat;
int pitch=devMat->pitch;
i = cublasSetMatrix(lda,nc,sizeof(float), (void*) val,lda, (void *)devMat->cM, pitch);
@ -252,14 +229,10 @@ int writeDnsDeviceFloat(void* deviceMat, float* val, int lda, int nc)
fprintf(stderr,"From routine : %s : %d \n","writeDnsDeviceFloat",i);
}
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int writeDnsDeviceDouble(void* deviceMat, double* val, int lda, int nc)
{ int i;
#ifdef HAVE_SPGPU
struct DnsDevice *devMat = (struct DnsDevice *) deviceMat;
int pitch=devMat->pitch;
i = cublasSetMatrix(lda,nc,sizeof(double), (void*) val,lda, (void *)devMat->cM, pitch);
@ -267,15 +240,11 @@ int writeDnsDeviceDouble(void* deviceMat, double* val, int lda, int nc)
fprintf(stderr,"From routine : %s : %d \n","writeDnsDeviceDouble",i);
}
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int writeDnsDeviceFloatComplex(void* deviceMat, float complex* val, int lda, int nc)
{ int i;
#ifdef HAVE_SPGPU
struct DnsDevice *devMat = (struct DnsDevice *) deviceMat;
int pitch=devMat->pitch;
i = cublasSetMatrix(lda,nc,sizeof(cuFloatComplex), (void*) val,lda, (void *)devMat->cM, pitch);
@ -283,14 +252,10 @@ int writeDnsDeviceFloatComplex(void* deviceMat, float complex* val, int lda, int
fprintf(stderr,"From routine : %s : %d \n","writeDnsDeviceFloatComplex",i);
}
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int writeDnsDeviceDoubleComplex(void* deviceMat, double complex* val, int lda, int nc)
{ int i;
#ifdef HAVE_SPGPU
struct DnsDevice *devMat = (struct DnsDevice *) deviceMat;
int pitch=devMat->pitch;
i = cublasSetMatrix(lda,nc,sizeof(cuDoubleComplex), (void*) val,lda, (void *)devMat->cM, pitch);
@ -298,15 +263,11 @@ int writeDnsDeviceDoubleComplex(void* deviceMat, double complex* val, int lda, i
fprintf(stderr,"From routine : %s : %d \n","writeDnsDeviceDoubleComplex",i);
}
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int readDnsDeviceFloat(void* deviceMat, float* val, int lda, int nc)
{ int i;
#ifdef HAVE_SPGPU
struct DnsDevice *devMat = (struct DnsDevice *) deviceMat;
int pitch=devMat->pitch;
i = cublasGetMatrix(lda,nc,sizeof(float), (void*) val,lda, (void *)devMat->cM, pitch);
@ -314,14 +275,10 @@ int readDnsDeviceFloat(void* deviceMat, float* val, int lda, int nc)
fprintf(stderr,"From routine : %s : %d \n","readDnsDeviceFloat",i);
}
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int readDnsDeviceDouble(void* deviceMat, double* val, int lda, int nc)
{ int i;
#ifdef HAVE_SPGPU
struct DnsDevice *devMat = (struct DnsDevice *) deviceMat;
int pitch=devMat->pitch;
i = cublasGetMatrix(lda,nc,sizeof(double), (void*) val,lda, (void *)devMat->cM, pitch);
@ -329,15 +286,11 @@ int readDnsDeviceDouble(void* deviceMat, double* val, int lda, int nc)
fprintf(stderr,"From routine : %s : %d \n","readDnsDeviceDouble",i);
}
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int readDnsDeviceFloatComplex(void* deviceMat, float complex* val, int lda, int nc)
{ int i;
#ifdef HAVE_SPGPU
struct DnsDevice *devMat = (struct DnsDevice *) deviceMat;
int pitch=devMat->pitch;
i = cublasGetMatrix(lda,nc,sizeof(cuFloatComplex), (void*) val,lda, (void *)devMat->cM, pitch);
@ -345,14 +298,10 @@ int readDnsDeviceFloatComplex(void* deviceMat, float complex* val, int lda, int
fprintf(stderr,"From routine : %s : %d \n","readDnsDeviceFloatComplex",i);
}
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int readDnsDeviceDoubleComplex(void* deviceMat, double complex* val, int lda, int nc)
{ int i;
#ifdef HAVE_SPGPU
struct DnsDevice *devMat = (struct DnsDevice *) deviceMat;
int pitch=devMat->pitch;
i = cublasGetMatrix(lda,nc,sizeof(cuDoubleComplex), (void*) val,lda, (void *)devMat->cM, pitch);
@ -360,24 +309,13 @@ int readDnsDeviceDoubleComplex(void* deviceMat, double complex* val, int lda, in
fprintf(stderr,"From routine : %s : %d \n","readDnsDeviceDoubleComplex",i);
}
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int getDnsDevicePitch(void* deviceMat)
{ int i;
struct DnsDevice *devMat = (struct DnsDevice *) deviceMat;
#ifdef HAVE_SPGPU
i = devMat->pitch;
return(i);
#else
return SPGPU_UNSUPPORTED;
#endif
}
#endif

5
cuda/dnsdev.h
Unescape Escape View File

@ -33,7 +33,6 @@
#ifndef _DNSDEV_H_
#define _DNSDEV_H_
#if defined(HAVE_SPGPU)
#include "cintrf.h"
#include "cuComplex.h"
#include "cublas_v2.h"
@ -115,8 +114,4 @@ int getDnsDevicePitch(void* deviceMat);
//int spmvDnsDeviceFloat(void *deviceMat, float* alpha, void* deviceX, float* beta, void* deviceY);
//int spmvDnsDeviceDouble(void *deviceMat, double* alpha, void* deviceX, double* beta, void* deviceY);
#else
#define CINTRF_UNSUPPORTED -1
#endif
#endif

5
cuda/dnsdev_mod.F90
Unescape Escape View File

@ -44,8 +44,6 @@ module dnsdev_mod
integer(c_int) :: firstIndex
end type dnsdev_parms
#ifdef HAVE_SPGPU
interface
function FgetDnsDeviceParams(rows, columns, elementType, firstIndex) &
& result(res) bind(c,name='getDnsDeviceParams')
@ -269,7 +267,4 @@ module dnsdev_mod
end interface
#endif
end module dnsdev_mod

4
cuda/dvectordev.c
Unescape Escape View File

@ -32,7 +32,6 @@
#include <stdio.h>
#include <complex.h>
#if defined(HAVE_SPGPU)
//#include "utils.h"
//#include "common.h"
#include "dvectordev.h"
@ -300,6 +299,3 @@ int absMultiVecDeviceDouble(int n, double alpha, void *deviceVecA)
return(i);
}
#endif

3
cuda/dvectordev.h
Unescape Escape View File

@ -31,7 +31,6 @@
#pragma once
#if defined(HAVE_SPGPU)
//#include "utils.h"
#include "vectordev.h"
#include "cuda_runtime.h"
@ -74,5 +73,3 @@ int axybzMultiVecDeviceDouble(int n, double alpha, void *deviceVecA,
int absMultiVecDeviceDouble(int n, double alpha, void *deviceVecA);
int absMultiVecDeviceDouble2(int n, double alpha, void *deviceVecA, void *deviceVecB);
#endif

87
cuda/elldev.c
Unescape Escape View File

@ -31,8 +31,6 @@
#include <sys/time.h>
#include "elldev.h"
#if defined(HAVE_SPGPU)
#define PASS_RS 0
EllDeviceParams getEllDeviceParams(unsigned int rows, unsigned int maxRowSize,
@ -140,7 +138,6 @@ int FallocEllDevice(void** deviceMat,unsigned int rows, unsigned int maxRowSize,
unsigned int columns, unsigned int elementType,
unsigned int firstIndex)
{ int i;
#ifdef HAVE_SPGPU
EllDeviceParams p;
p = getEllDeviceParams(rows, maxRowSize, nnzeros, columns, elementType, firstIndex);
@ -149,9 +146,6 @@ int FallocEllDevice(void** deviceMat,unsigned int rows, unsigned int maxRowSize,
fprintf(stderr,"From routine : %s : %d \n","FallocEllDevice",i);
}
return(i);
#else
return SPGPU_UNSUPPORTED;
#endif
}
void sspmdmm_gpu(float *z,int s, int vPitch, float *y, float alpha, float* cM, int* rP, int* rS,
@ -182,7 +176,6 @@ int spmvEllDeviceFloat(void *deviceMat, float alpha, void* deviceX,
struct MultiVectDevice *x = (struct MultiVectDevice *) deviceX;
struct MultiVectDevice *y = (struct MultiVectDevice *) deviceY;
#ifdef HAVE_SPGPU
#ifdef VERBOSE
__assert(x->count_ == x->count_, "ERROR: x and y don't share the same number of vectors");
__assert(x->size_ >= devMat->columns, "ERROR: x vector's size is not >= to matrix size (columns)");
@ -196,9 +189,6 @@ int spmvEllDeviceFloat(void *deviceMat, float alpha, void* deviceX,
devMat->avgRowSize, devMat->maxRowSize, devMat->rows, devMat->pitch,
(float *)x->v_, beta, devMat->baseIndex);
return(i);
#else
return SPGPU_UNSUPPORTED;
#endif
}
@ -234,7 +224,6 @@ int spmvEllDeviceDouble(void *deviceMat, double alpha, void* deviceX,
struct MultiVectDevice *x = (struct MultiVectDevice *) deviceX;
struct MultiVectDevice *y = (struct MultiVectDevice *) deviceY;
#ifdef HAVE_SPGPU
/*spgpuDellspmv (handle, (double*) y->v_, (double*)y->v_, alpha, (double*) devMat->cM, devMat->rP, devMat->cMPitch, devMat->rPPitch, devMat->rS, devMat->rows, (double*)x->v_, beta, devMat->baseIndex);*/
/* fprintf(stderr,"From spmvEllDouble: mat %d %d %d %d y %d %d \n", */
/* devMat->avgRowSize, devMat->maxRowSize, devMat->rows, */
@ -246,9 +235,6 @@ int spmvEllDeviceDouble(void *deviceMat, double alpha, void* deviceX,
(double *)x->v_, beta, devMat->baseIndex);
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
void
@ -281,7 +267,6 @@ int spmvEllDeviceFloatComplex(void *deviceMat, float complex alpha, void* device
struct MultiVectDevice *x = (struct MultiVectDevice *) deviceX;
struct MultiVectDevice *y = (struct MultiVectDevice *) deviceY;
#ifdef HAVE_SPGPU
cuFloatComplex a = make_cuFloatComplex(crealf(alpha),cimagf(alpha));
cuFloatComplex b = make_cuFloatComplex(crealf(beta),cimagf(beta));
cspmdmm_gpu ((cuFloatComplex *)y->v_, y->count_, y->pitch_, (cuFloatComplex *)y->v_, a, (cuFloatComplex *)devMat->cM,
@ -289,9 +274,6 @@ int spmvEllDeviceFloatComplex(void *deviceMat, float complex alpha, void* device
(cuFloatComplex *)x->v_, b, devMat->baseIndex);
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
void
@ -323,7 +305,6 @@ int spmvEllDeviceDoubleComplex(void *deviceMat, double complex alpha, void* devi
struct MultiVectDevice *x = (struct MultiVectDevice *) deviceX;
struct MultiVectDevice *y = (struct MultiVectDevice *) deviceY;
#ifdef HAVE_SPGPU
cuDoubleComplex a = make_cuDoubleComplex(creal(alpha),cimag(alpha));
cuDoubleComplex b = make_cuDoubleComplex(creal(beta),cimag(beta));
zspmdmm_gpu ((cuDoubleComplex *)y->v_, y->count_, y->pitch_, (cuDoubleComplex *)y->v_, a, (cuDoubleComplex *)devMat->cM,
@ -331,14 +312,10 @@ int spmvEllDeviceDoubleComplex(void *deviceMat, double complex alpha, void* devi
devMat->pitch, (cuDoubleComplex *)x->v_, b, devMat->baseIndex);
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int writeEllDeviceFloat(void* deviceMat, float* val, int* ja, int ldj, int* irn, int *idiag)
{ int i;
#ifdef HAVE_SPGPU
struct EllDevice *devMat = (struct EllDevice *) deviceMat;
// Ex updateFromHost function
i = writeRemoteBuffer((void*) val, (void *)devMat->cM, devMat->allocsize*sizeof(float));
@ -350,14 +327,10 @@ int writeEllDeviceFloat(void* deviceMat, float* val, int* ja, int ldj, int* irn,
fprintf(stderr,"From routine : %s : %d \n","writeEllDeviceFloat",i);
}*/
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int writeEllDeviceDouble(void* deviceMat, double* val, int* ja, int ldj, int* irn, int *idiag)
{ int i;
#ifdef HAVE_SPGPU
struct EllDevice *devMat = (struct EllDevice *) deviceMat;
// Ex updateFromHost function
i = writeRemoteBuffer((void*) val, (void *)devMat->cM, devMat->allocsize*sizeof(double));
@ -370,14 +343,10 @@ int writeEllDeviceDouble(void* deviceMat, double* val, int* ja, int ldj, int* ir
fprintf(stderr,"From routine : %s : %d \n","writeEllDeviceDouble",i);
}
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int writeEllDeviceFloatComplex(void* deviceMat, float complex* val, int* ja, int ldj, int* irn, int *idiag)
{ int i;
#ifdef HAVE_SPGPU
struct EllDevice *devMat = (struct EllDevice *) deviceMat;
// Ex updateFromHost function
i = writeRemoteBuffer((void*) val, (void *)devMat->cM, devMat->allocsize*sizeof(cuFloatComplex));
@ -390,14 +359,10 @@ int writeEllDeviceFloatComplex(void* deviceMat, float complex* val, int* ja, int
fprintf(stderr,"From routine : %s : %d \n","writeEllDeviceDouble",i);
}*/
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int writeEllDeviceDoubleComplex(void* deviceMat, double complex* val, int* ja, int ldj, int* irn, int *idiag)
{ int i;
#ifdef HAVE_SPGPU
struct EllDevice *devMat = (struct EllDevice *) deviceMat;
// Ex updateFromHost function
i = writeRemoteBuffer((void*) val, (void *)devMat->cM, devMat->allocsize*sizeof(cuDoubleComplex));
@ -410,14 +375,10 @@ int writeEllDeviceDoubleComplex(void* deviceMat, double complex* val, int* ja, i
fprintf(stderr,"From routine : %s : %d \n","writeEllDeviceDouble",i);
}*/
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int readEllDeviceFloat(void* deviceMat, float* val, int* ja, int ldj, int* irn, int *idiag)
{ int i;
#ifdef HAVE_SPGPU
struct EllDevice *devMat = (struct EllDevice *) deviceMat;
i = readRemoteBuffer((void *) val, (void *)devMat->cM, devMat->allocsize*sizeof(float));
i = readRemoteBuffer((void *) ja, (void *)devMat->rP, devMat->allocsize*sizeof(int));
@ -428,14 +389,10 @@ int readEllDeviceFloat(void* deviceMat, float* val, int* ja, int ldj, int* irn,
fprintf(stderr,"From routine : %s : %d \n","readEllDeviceFloat",i);
}*/
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int readEllDeviceDouble(void* deviceMat, double* val, int* ja, int ldj, int* irn, int *idiag)
{ int i;
#ifdef HAVE_SPGPU
struct EllDevice *devMat = (struct EllDevice *) deviceMat;
i = readRemoteBuffer((void *) val, (void *)devMat->cM, devMat->allocsize*sizeof(double));
i = readRemoteBuffer((void *) ja, (void *)devMat->rP, devMat->allocsize*sizeof(int));
@ -445,14 +402,10 @@ int readEllDeviceDouble(void* deviceMat, double* val, int* ja, int ldj, int* irn
fprintf(stderr,"From routine : %s : %d \n","readEllDeviceDouble",i);
}*/
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int readEllDeviceFloatComplex(void* deviceMat, float complex* val, int* ja, int ldj, int* irn, int *idiag)
{ int i;
#ifdef HAVE_SPGPU
struct EllDevice *devMat = (struct EllDevice *) deviceMat;
i = readRemoteBuffer((void *) val, (void *)devMat->cM, devMat->allocsize*sizeof(cuFloatComplex));
i = readRemoteBuffer((void *) ja, (void *)devMat->rP, devMat->allocsize*sizeof(int));
@ -462,14 +415,10 @@ int readEllDeviceFloatComplex(void* deviceMat, float complex* val, int* ja, int
fprintf(stderr,"From routine : %s : %d \n","readEllDeviceDouble",i);
}*/
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int readEllDeviceDoubleComplex(void* deviceMat, double complex* val, int* ja, int ldj, int* irn, int *idiag)
{ int i;
#ifdef HAVE_SPGPU
struct EllDevice *devMat = (struct EllDevice *) deviceMat;
i = readRemoteBuffer((void *) val, (void *)devMat->cM, devMat->allocsize*sizeof(cuDoubleComplex));
i = readRemoteBuffer((void *) ja, (void *)devMat->rP, devMat->allocsize*sizeof(int));
@ -479,32 +428,21 @@ int readEllDeviceDoubleComplex(void* deviceMat, double complex* val, int* ja, in
fprintf(stderr,"From routine : %s : %d \n","readEllDeviceDouble",i);
}*/
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int getEllDevicePitch(void* deviceMat)
{ int i;
struct EllDevice *devMat = (struct EllDevice *) deviceMat;
#ifdef HAVE_SPGPU
i = devMat->pitch; //old
//i = getPitchEllDevice(deviceMat);
return(i);
#else
return SPGPU_UNSUPPORTED;
#endif
}
int getEllDeviceMaxRowSize(void* deviceMat)
{ int i;
struct EllDevice *devMat = (struct EllDevice *) deviceMat;
#ifdef HAVE_SPGPU
i = devMat->maxRowSize;
return(i);
#else
return SPGPU_UNSUPPORTED;
#endif
}
@ -515,7 +453,6 @@ int getEllDeviceMaxRowSize(void* deviceMat)
int psiCopyCooToElgFloat(int nr, int nc, int nza, int hacksz, int ldv, int nzm, int *irn,
int *idisp, int *ja, float *val, void *deviceMat)
{ int i;
#ifdef HAVE_SPGPU
struct EllDevice *devMat = (struct EllDevice *) deviceMat;
float *devVal;
int *devIdisp, *devJa;
@ -548,9 +485,6 @@ int psiCopyCooToElgFloat(int nr, int nc, int nza, int hacksz, int ldv, int nzm,
fprintf(stderr,"From routine : %s : %d \n","writeEllDeviceFloat",i);
}
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
@ -558,7 +492,6 @@ int psiCopyCooToElgFloat(int nr, int nc, int nza, int hacksz, int ldv, int nzm,
int psiCopyCooToElgDouble(int nr, int nc, int nza, int hacksz, int ldv, int nzm, int *irn,
int *idisp, int *ja, double *val, void *deviceMat)
{ int i;
#ifdef HAVE_SPGPU
struct EllDevice *devMat = (struct EllDevice *) deviceMat;
double *devVal;
int *devIdisp, *devJa;
@ -591,16 +524,12 @@ int psiCopyCooToElgDouble(int nr, int nc, int nza, int hacksz, int ldv, int nzm,
fprintf(stderr,"From routine : %s : %d \n","writeEllDeviceDouble",i);
}
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int psiCopyCooToElgFloatComplex(int nr, int nc, int nza, int hacksz, int ldv, int nzm, int *irn,
int *idisp, int *ja, float complex *val, void *deviceMat)
{ int i;
#ifdef HAVE_SPGPU
struct EllDevice *devMat = (struct EllDevice *) deviceMat;
float complex *devVal;
int *devIdisp, *devJa;
@ -633,9 +562,6 @@ int psiCopyCooToElgFloatComplex(int nr, int nc, int nza, int hacksz, int ldv, in
fprintf(stderr,"From routine : %s : %d \n","writeEllDeviceFloatComplex",i);
}
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
@ -643,7 +569,6 @@ int psiCopyCooToElgFloatComplex(int nr, int nc, int nza, int hacksz, int ldv, in
int psiCopyCooToElgDoubleComplex(int nr, int nc, int nza, int hacksz, int ldv, int nzm, int *irn,
int *idisp, int *ja, double complex *val, void *deviceMat)
{ int i;
#ifdef HAVE_SPGPU
struct EllDevice *devMat = (struct EllDevice *) deviceMat;
double complex *devVal;
int *devIdisp, *devJa;
@ -676,15 +601,11 @@ int psiCopyCooToElgDoubleComplex(int nr, int nc, int nza, int hacksz, int ldv, i
fprintf(stderr,"From routine : %s : %d \n","writeEllDeviceDoubleComplex",i);
}
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int dev_csputEllDeviceFloat(void* deviceMat, int nnz, void *ia, void *ja, void *val)
{ int i;
#ifdef HAVE_SPGPU
struct EllDevice *devMat = (struct EllDevice *) deviceMat;
struct MultiVectDevice *devVal = (struct MultiVectDevice *) val;
struct MultiVectDevice *devIa = (struct MultiVectDevice *) ia;
@ -699,13 +620,11 @@ int dev_csputEllDeviceFloat(void* deviceMat, int nnz, void *ia, void *ja, void *
devMat->rP,devMat->pitch, devMat->pitch, devMat->rS,
nnz, devIa->v_, devJa->v_, (float *) devVal->v_, 1);
#endif
return SPGPU_SUCCESS;
}
int dev_csputEllDeviceDouble(void* deviceMat, int nnz, void *ia, void *ja, void *val)
{ int i;
#ifdef HAVE_SPGPU
struct EllDevice *devMat = (struct EllDevice *) deviceMat;
struct MultiVectDevice *devVal = (struct MultiVectDevice *) val;
struct MultiVectDevice *devIa = (struct MultiVectDevice *) ia;
@ -720,7 +639,6 @@ int dev_csputEllDeviceDouble(void* deviceMat, int nnz, void *ia, void *ja, void
devMat->rP,devMat->pitch, devMat->pitch, devMat->rS,
nnz, devIa->v_, devJa->v_, (double *) devVal->v_, 1);
#endif
return SPGPU_SUCCESS;
}
@ -728,7 +646,6 @@ int dev_csputEllDeviceDouble(void* deviceMat, int nnz, void *ia, void *ja, void
int dev_csputEllDeviceFloatComplex(void* deviceMat, int nnz,
void *ia, void *ja, void *val)
{ int i;
#ifdef HAVE_SPGPU
struct EllDevice *devMat = (struct EllDevice *) deviceMat;
struct MultiVectDevice *devVal = (struct MultiVectDevice *) val;
struct MultiVectDevice *devIa = (struct MultiVectDevice *) ia;
@ -743,14 +660,12 @@ int dev_csputEllDeviceFloatComplex(void* deviceMat, int nnz,
devMat->rP,devMat->pitch, devMat->pitch, devMat->rS,
nnz, devIa->v_, devJa->v_, (cuFloatComplex *) devVal->v_, 1);
#endif
return SPGPU_SUCCESS;
}
int dev_csputEllDeviceDoubleComplex(void* deviceMat, int nnz,
void *ia, void *ja, void *val)
{ int i;
#ifdef HAVE_SPGPU
struct EllDevice *devMat = (struct EllDevice *) deviceMat;
struct MultiVectDevice *devVal = (struct MultiVectDevice *) val;
struct MultiVectDevice *devIa = (struct MultiVectDevice *) ia;
@ -765,9 +680,7 @@ int dev_csputEllDeviceDoubleComplex(void* deviceMat, int nnz,
devMat->rP,devMat->pitch, devMat->pitch, devMat->rS,
nnz, devIa->v_, devJa->v_, (cuDoubleComplex *) devVal->v_, 1);
#endif
return SPGPU_SUCCESS;
}
#endif

6
cuda/elldev.h
Unescape Escape View File

@ -33,12 +33,10 @@
#ifndef _ELLDEV_H_
#define _ELLDEV_H_
#if defined(HAVE_SPGPU)
#include "cintrf.h"
#include "cuComplex.h"
#include "ell.h"
struct EllDevice
{
// Compressed matrix
@ -176,8 +174,4 @@ int getEllDevicePitch(void* deviceMat);
//int spmvEllDeviceFloat(void *deviceMat, float* alpha, void* deviceX, float* beta, void* deviceY);
//int spmvEllDeviceDouble(void *deviceMat, double* alpha, void* deviceX, double* beta, void* deviceY);
#else
#define CINTRF_UNSUPPORTED -1
#endif
#endif

5
cuda/elldev_mod.F90
Unescape Escape View File

@ -44,8 +44,6 @@ module elldev_mod
integer(c_int) :: firstIndex
end type elldev_parms
#ifdef HAVE_SPGPU
interface
function FgetEllDeviceParams(rows, maxRowSize, nnzeros, columns, elementType, firstIndex) &
& result(res) bind(c,name='getEllDeviceParams')
@ -320,7 +318,4 @@ module elldev_mod
end function spmvEllDeviceDoubleComplex
end interface
#endif
end module elldev_mod

4
cuda/fcusparse.c
Unescape Escape View File

@ -33,7 +33,6 @@
#include <stdio.h>
#include <stdlib.h>
#ifdef HAVE_SPGPU
#include <cuda_runtime.h>
#include "cintrf.h"
#include "fcusparse.h"
@ -72,6 +71,3 @@ cusparseHandle_t *getHandle()
return(cusparse_handle);
}
#endif

2
cuda/fcusparse.h
Unescape Escape View File

@ -33,7 +33,6 @@
#ifndef FCUSPARSE_
#define FCUSPARSE_
#ifdef HAVE_SPGPU
#include <cuda_runtime.h>
#if CUDA_SHORT_VERSION <= 10
#include <cusparse_v2.h>
@ -67,4 +66,3 @@ cusparseHandle_t *getHandle();
}
#endif
#endif

39
cuda/hdiagdev.c
Unescape Escape View File

@ -34,10 +34,7 @@
#include <sys/stat.h>
#include <fcntl.h>
#include <unistd.h>
#if defined(HAVE_SPGPU)
#define DEBUG 0
void freeHdiagDevice(void* remoteMatrix)
{
struct HdiagDevice *devMat = (struct HdiagDevice *) remoteMatrix;
@ -138,7 +135,6 @@ int FallocHdiagDevice(void** deviceMat, unsigned int rows, unsigned int cols,
unsigned int allocationHeight, unsigned int hackSize,
unsigned int hackCount, unsigned int elementType)
{ int i=0;
#ifdef HAVE_SPGPU
HdiagDeviceParams p;
p = getHdiagDeviceParams(rows, cols, allocationHeight, hackSize, hackCount,elementType);
@ -152,17 +148,12 @@ int FallocHdiagDevice(void** deviceMat, unsigned int rows, unsigned int cols,
fprintf(stderr,"From routine : %s : %d \n","FallocEllDevice",i);
}
return(i);
#else
return SPGPU_UNSUPPORTED;
#endif
}
int writeHdiagDeviceDouble(void* deviceMat, double* val, int* hdiaOffsets, int *hackOffsets)
{ int i=0,fo,fa,j,k,p;
char buf_a[255], buf_o[255],tmp[255];
#ifdef HAVE_SPGPU
struct HdiagDevice *devMat = (struct HdiagDevice *) deviceMat;
i=SPGPU_SUCCESS;
@ -216,9 +207,6 @@ int writeHdiagDeviceDouble(void* deviceMat, double* val, int* hdiaOffsets, int *
return SPGPU_SUCCESS;
else
return SPGPU_UNSUPPORTED;
#else
return SPGPU_UNSUPPORTED;
#endif
}
@ -227,15 +215,12 @@ long long int sizeofHdiagDeviceDouble(void* deviceMat)
{ int i=0,fo,fa;
int *hoff=NULL,*hackoff=NULL;
long long int memsize=0;
#ifdef HAVE_SPGPU
struct HdiagDevice *devMat = (struct HdiagDevice *) deviceMat;
memsize += (devMat->hackCount+1)*sizeof(int);
memsize += devMat->allocationHeight*sizeof(int);
memsize += devMat->allocationHeight*devMat->hackSize*sizeof(double);
#endif
return(memsize);
}
@ -243,7 +228,6 @@ long long int sizeofHdiagDeviceDouble(void* deviceMat)
int readHdiagDeviceDouble(void* deviceMat, double* a, int* off)
{ int i;
#ifdef HAVE_SPGPU
struct HdiagDevice *devMat = (struct HdiagDevice *) deviceMat;
/* i = readRemoteBuffer((void *) a, (void *)devMat->cM,devMat->rows*devMat->diags*sizeof(double)); */
/* i = readRemoteBuffer((void *) off, (void *)devMat->off, devMat->diags*sizeof(int)); */
@ -253,9 +237,6 @@ int readHdiagDeviceDouble(void* deviceMat, double* a, int* off)
fprintf(stderr,"From routine : %s : %d \n","readEllDeviceDouble",i);
}*/
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int spmvHdiagDeviceDouble(void *deviceMat, double alpha, void* deviceX,
@ -266,7 +247,6 @@ int spmvHdiagDeviceDouble(void *deviceMat, double alpha, void* deviceX,
struct MultiVectDevice *y = (struct MultiVectDevice *) deviceY;
spgpuHandle_t handle=psb_cudaGetHandle();
#ifdef HAVE_SPGPU
#ifdef VERBOSE
/*__assert(x->count_ == x->count_, "ERROR: x and y don't share the same number of vectors");*/
/*__assert(x->size_ >= devMat->columns, "ERROR: x vector's size is not >= to matrix size (columns)");*/
@ -285,15 +265,11 @@ int spmvHdiagDeviceDouble(void *deviceMat, double alpha, void* deviceX,
//cudaSync();
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int writeHdiagDeviceFloat(void* deviceMat, float* val, int* hdiaOffsets, int *hackOffsets)
{ int i=0,fo,fa,j,k,p;
char buf_a[255], buf_o[255],tmp[255];
#ifdef HAVE_SPGPU
struct HdiagDevice *devMat = (struct HdiagDevice *) deviceMat;
i=SPGPU_SUCCESS;
@ -347,9 +323,6 @@ int writeHdiagDeviceFloat(void* deviceMat, float* val, int* hdiaOffsets, int *ha
return SPGPU_SUCCESS;
else
return SPGPU_UNSUPPORTED;
#else
return SPGPU_UNSUPPORTED;
#endif
}
@ -358,7 +331,6 @@ long long int sizeofHdiagDeviceFloat(void* deviceMat)
{ int i=0,fo,fa;
int *hoff=NULL,*hackoff=NULL;
long long int memsize=0;
#ifdef HAVE_SPGPU
struct HdiagDevice *devMat = (struct HdiagDevice *) deviceMat;
@ -366,7 +338,6 @@ long long int sizeofHdiagDeviceFloat(void* deviceMat)
memsize += devMat->allocationHeight*sizeof(int);
memsize += devMat->allocationHeight*devMat->hackSize*sizeof(float);
#endif
return(memsize);
}
@ -374,7 +345,6 @@ long long int sizeofHdiagDeviceFloat(void* deviceMat)
int readHdiagDeviceFloat(void* deviceMat, float* a, int* off)
{ int i;
#ifdef HAVE_SPGPU
struct HdiagDevice *devMat = (struct HdiagDevice *) deviceMat;
/* i = readRemoteBuffer((void *) a, (void *)devMat->cM,devMat->rows*devMat->diags*sizeof(float)); */
/* i = readRemoteBuffer((void *) off, (void *)devMat->off, devMat->diags*sizeof(int)); */
@ -384,9 +354,6 @@ int readHdiagDeviceFloat(void* deviceMat, float* a, int* off)
fprintf(stderr,"From routine : %s : %d \n","readEllDeviceFloat",i);
}*/
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int spmvHdiagDeviceFloat(void *deviceMat, float alpha, void* deviceX,
@ -397,7 +364,6 @@ int spmvHdiagDeviceFloat(void *deviceMat, float alpha, void* deviceX,
struct MultiVectDevice *y = (struct MultiVectDevice *) deviceY;
spgpuHandle_t handle=psb_cudaGetHandle();
#ifdef HAVE_SPGPU
#ifdef VERBOSE
/*__assert(x->count_ == x->count_, "ERROR: x and y don't share the same number of vectors");*/
/*__assert(x->size_ >= devMat->columns, "ERROR: x vector's size is not >= to matrix size (columns)");*/
@ -416,10 +382,5 @@ int spmvHdiagDeviceFloat(void *deviceMat, float alpha, void* deviceX,
//cudaSync();
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
#endif

5
cuda/hdiagdev.h
Unescape Escape View File

@ -32,7 +32,6 @@
#ifndef _HDIAGDEV_H_
#define _HDIAGDEV_H_
#ifdef HAVE_SPGPU
#include "cintrf.h"
#include "hdia.h"
@ -104,8 +103,4 @@ int spmvHdiagDeviceDouble(void *deviceMat, double alpha, void* deviceX,
double beta, void* deviceY);
#else
#define CINTRF_UNSUPPORTED -1
#endif
#endif

4
cuda/hdiagdev_mod.F90
Unescape Escape View File

@ -43,8 +43,6 @@ module hdiagdev_mod
integer(c_int) :: allocationHeight
end type hdiagdev_parms
#ifdef HAVE_SPGPU
! interface computeHdiaHacksCount
! function computeHdiaHacksCountDouble(allocationHeight,hackOffsets,hackSize, &
! & diaValues,diaValuesPitch,diags,rows)&
@ -198,6 +196,4 @@ module hdiagdev_mod
!!$ end function spmvHdiagDeviceDoubleComplex
end interface spmvHdiagDevice
#endif
end module hdiagdev_mod

75
cuda/hlldev.c
Unescape Escape View File

@ -30,7 +30,6 @@
#include "hlldev.h"
#if defined(HAVE_SPGPU)
//new
HllDeviceParams bldHllDeviceParams(unsigned int hksize, unsigned int rows, unsigned int nzeros,
unsigned int allocsize, unsigned int elementType, unsigned int firstIndex)
@ -147,7 +146,6 @@ int FallocHllDevice(void** deviceMat,unsigned int hksize, unsigned int rows, un
unsigned int allocsize,
unsigned int elementType, unsigned int firstIndex)
{ int i;
#ifdef HAVE_SPGPU
HllDeviceParams p;
p = bldHllDeviceParams(hksize, rows, nzeros, allocsize, elementType, firstIndex);
@ -156,9 +154,6 @@ int FallocHllDevice(void** deviceMat,unsigned int hksize, unsigned int rows, un
fprintf(stderr,"From routine : %s : %d \n","FallocEllDevice",i);
}
return(i);
#else
return SPGPU_UNSUPPORTED;
#endif
}
@ -170,7 +165,6 @@ int spmvHllDeviceFloat(void *deviceMat, float alpha, void* deviceX,
struct MultiVectDevice *y = (struct MultiVectDevice *) deviceY;
spgpuHandle_t handle=psb_cudaGetHandle();
#ifdef HAVE_SPGPU
#ifdef VERBOSE
/*__assert(x->count_ == x->count_, "ERROR: x and y don't share the same number of vectors");*/
/*__assert(x->size_ >= devMat->columns, "ERROR: x vector's size is not >= to matrix size (columns)");*/
@ -185,9 +179,6 @@ int spmvHllDeviceFloat(void *deviceMat, float alpha, void* deviceX,
devMat->avgNzr, devMat->rows, (float *)x->v_, beta, devMat->baseIndex);
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
//new
@ -199,7 +190,6 @@ int spmvHllDeviceDouble(void *deviceMat, double alpha, void* deviceX,
struct MultiVectDevice *y = (struct MultiVectDevice *) deviceY;
spgpuHandle_t handle=psb_cudaGetHandle();
#ifdef HAVE_SPGPU
#ifdef VERBOSE
/*__assert(x->count_ == x->count_, "ERROR: x and y don't share the same number of vectors");*/
/*__assert(x->size_ >= devMat->columns, "ERROR: x vector's size is not >= to matrix size (columns)");*/
@ -214,9 +204,6 @@ int spmvHllDeviceDouble(void *deviceMat, double alpha, void* deviceX,
devMat->avgNzr, devMat->rows, (double *)x->v_, beta, devMat->baseIndex);
//cudaSync();
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int spmvHllDeviceFloatComplex(void *deviceMat, float complex alpha, void* deviceX,
@ -227,7 +214,6 @@ int spmvHllDeviceFloatComplex(void *deviceMat, float complex alpha, void* device
struct MultiVectDevice *y = (struct MultiVectDevice *) deviceY;
spgpuHandle_t handle=psb_cudaGetHandle();
#ifdef HAVE_SPGPU
cuFloatComplex a = make_cuFloatComplex(crealf(alpha),cimagf(alpha));
cuFloatComplex b = make_cuFloatComplex(crealf(beta),cimagf(beta));
#ifdef VERBOSE
@ -244,9 +230,6 @@ int spmvHllDeviceFloatComplex(void *deviceMat, float complex alpha, void* device
devMat->avgNzr, devMat->rows, (cuFloatComplex *)x->v_, b, devMat->baseIndex);
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int spmvHllDeviceDoubleComplex(void *deviceMat, double complex alpha, void* deviceX,
@ -257,7 +240,6 @@ int spmvHllDeviceDoubleComplex(void *deviceMat, double complex alpha, void* devi
struct MultiVectDevice *y = (struct MultiVectDevice *) deviceY;
spgpuHandle_t handle=psb_cudaGetHandle();
#ifdef HAVE_SPGPU
cuDoubleComplex a = make_cuDoubleComplex(creal(alpha),cimag(alpha));
cuDoubleComplex b = make_cuDoubleComplex(creal(beta),cimag(beta));
#ifdef VERBOSE
@ -271,14 +253,10 @@ int spmvHllDeviceDoubleComplex(void *deviceMat, double complex alpha, void* devi
devMat->avgNzr,devMat->rows, (cuDoubleComplex *)x->v_, b, devMat->baseIndex);
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int writeHllDeviceFloat(void* deviceMat, float* val, int* ja, int *hkoffs, int* irn, int *idiag)
{ int i;
#ifdef HAVE_SPGPU
HllDevice *devMat = (HllDevice *) deviceMat;
// Ex updateFromHost function
i = writeRemoteBuffer((void*) val, (void *)devMat->cM, devMat->allocsize*sizeof(float));
@ -291,14 +269,10 @@ int writeHllDeviceFloat(void* deviceMat, float* val, int* ja, int *hkoffs, int*
fprintf(stderr,"From routine : %s : %d \n","writeEllDeviceFloat",i);
}*/
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int writeHllDeviceDouble(void* deviceMat, double* val, int* ja, int *hkoffs, int* irn, int *idiag)
{ int i;
#ifdef HAVE_SPGPU
HllDevice *devMat = (HllDevice *) deviceMat;
// Ex updateFromHost function
i = writeRemoteBuffer((void*) val, (void *)devMat->cM, devMat->allocsize*sizeof(double));
@ -311,14 +285,10 @@ int writeHllDeviceDouble(void* deviceMat, double* val, int* ja, int *hkoffs, int
fprintf(stderr,"From routine : %s : %d \n","writeEllDeviceDouble",i);
}*/
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int writeHllDeviceFloatComplex(void* deviceMat, float complex* val, int* ja, int *hkoffs, int* irn, int *idiag)
{ int i;
#ifdef HAVE_SPGPU
HllDevice *devMat = (HllDevice *) deviceMat;
// Ex updateFromHost function
i = writeRemoteBuffer((void*) val, (void *)devMat->cM, devMat->allocsize*sizeof(cuFloatComplex));
@ -331,14 +301,10 @@ int writeHllDeviceFloatComplex(void* deviceMat, float complex* val, int* ja, int
fprintf(stderr,"From routine : %s : %d \n","writeEllDeviceDouble",i);
}*/
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int writeHllDeviceDoubleComplex(void* deviceMat, double complex* val, int* ja, int *hkoffs, int* irn, int *idiag)
{ int i;
#ifdef HAVE_SPGPU
HllDevice *devMat = (HllDevice *) deviceMat;
// Ex updateFromHost function
i = writeRemoteBuffer((void*) val, (void *)devMat->cM, devMat->allocsize*sizeof(cuDoubleComplex));
@ -351,14 +317,10 @@ int writeHllDeviceDoubleComplex(void* deviceMat, double complex* val, int* ja, i
fprintf(stderr,"From routine : %s : %d \n","writeEllDeviceDouble",i);
}*/
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int readHllDeviceFloat(void* deviceMat, float* val, int* ja, int *hkoffs, int* irn, int *idiag)
{ int i;
#ifdef HAVE_SPGPU
HllDevice *devMat = (HllDevice *) deviceMat;
i = readRemoteBuffer((void *) val, (void *)devMat->cM, devMat->allocsize*sizeof(float));
i = readRemoteBuffer((void *) ja, (void *)devMat->rP, devMat->allocsize*sizeof(int));
@ -370,14 +332,10 @@ int readHllDeviceFloat(void* deviceMat, float* val, int* ja, int *hkoffs, int* i
fprintf(stderr,"From routine : %s : %d \n","readEllDeviceFloat",i);
}*/
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int readHllDeviceDouble(void* deviceMat, double* val, int* ja, int *hkoffs, int* irn, int *idiag)
{ int i;
#ifdef HAVE_SPGPU
HllDevice *devMat = (HllDevice *) deviceMat;
i = readRemoteBuffer((void *) val, (void *)devMat->cM, devMat->allocsize*sizeof(double));
i = readRemoteBuffer((void *) ja, (void *)devMat->rP, devMat->allocsize*sizeof(int));
@ -388,14 +346,10 @@ int readHllDeviceDouble(void* deviceMat, double* val, int* ja, int *hkoffs, int*
fprintf(stderr,"From routine : %s : %d \n","readEllDeviceDouble",i);
}*/
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int readHllDeviceFloatComplex(void* deviceMat, float complex* val, int* ja, int *hkoffs, int* irn, int *idiag)
{ int i;
#ifdef HAVE_SPGPU
HllDevice *devMat = (HllDevice *) deviceMat;
i = readRemoteBuffer((void *) val, (void *)devMat->cM, devMat->allocsize*sizeof(cuFloatComplex));
i = readRemoteBuffer((void *) ja, (void *)devMat->rP, devMat->allocsize*sizeof(int));
@ -406,14 +360,10 @@ int readHllDeviceFloatComplex(void* deviceMat, float complex* val, int* ja, int
fprintf(stderr,"From routine : %s : %d \n","readEllDeviceDouble",i);
}*/
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int readHllDeviceDoubleComplex(void* deviceMat, double complex* val, int* ja, int *hkoffs, int* irn, int *idiag)
{ int i;
#ifdef HAVE_SPGPU
HllDevice *devMat = (HllDevice *) deviceMat;
i = readRemoteBuffer((void *) val, (void *)devMat->cM, devMat->allocsize*sizeof(cuDoubleComplex));
i = readRemoteBuffer((void *) ja, (void *)devMat->rP, devMat->allocsize*sizeof(int));
@ -424,9 +374,6 @@ int readHllDeviceDoubleComplex(void* deviceMat, double complex* val, int* ja, in
fprintf(stderr,"From routine : %s : %d \n","readEllDeviceDouble",i);
}*/
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
// New copy routines.
@ -435,7 +382,6 @@ int psiCopyCooToHlgFloat(int nr, int nc, int nza, int hacksz, int noffs, int isz
int *irn, int *hoffs, int *idisp, int *ja,
float *val, void *deviceMat)
{ int i,j;
#ifdef HAVE_SPGPU
spgpuHandle_t handle;
HllDevice *devMat = (HllDevice *) deviceMat;
float *devVal;
@ -469,16 +415,12 @@ int psiCopyCooToHlgFloat(int nr, int nc, int nza, int hacksz, int noffs, int isz
fprintf(stderr,"From routine : %s : %d \n","writeHllDeviceFloat",i);
}
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int psiCopyCooToHlgDouble(int nr, int nc, int nza, int hacksz, int noffs, int isz,
int *irn, int *hoffs, int *idisp, int *ja,
double *val, void *deviceMat)
{ int i,j;
#ifdef HAVE_SPGPU
spgpuHandle_t handle;
HllDevice *devMat = (HllDevice *) deviceMat;
double *devVal;
@ -517,16 +459,12 @@ int psiCopyCooToHlgDouble(int nr, int nc, int nza, int hacksz, int noffs, int is
fprintf(stderr,"From routine : %s : %d \n","writeHllDeviceDouble",i);
}
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int psiCopyCooToHlgFloatComplex(int nr, int nc, int nza, int hacksz, int noffs, int isz,
int *irn, int *hoffs, int *idisp, int *ja,
float complex *val, void *deviceMat)
{ int i,j;
#ifdef HAVE_SPGPU
spgpuHandle_t handle;
HllDevice *devMat = (HllDevice *) deviceMat;
float complex *devVal;
@ -560,16 +498,12 @@ int psiCopyCooToHlgFloatComplex(int nr, int nc, int nza, int hacksz, int noffs,
fprintf(stderr,"From routine : %s : %d \n","writeHllDeviceFloatComplex",i);
}
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
int psiCopyCooToHlgDoubleComplex(int nr, int nc, int nza, int hacksz, int noffs, int isz,
int *irn, int *hoffs, int *idisp, int *ja,
double complex *val, void *deviceMat)
{ int i,j;
#ifdef HAVE_SPGPU
spgpuHandle_t handle;
HllDevice *devMat = (HllDevice *) deviceMat;
double complex *devVal;
@ -603,13 +537,4 @@ int psiCopyCooToHlgDoubleComplex(int nr, int nc, int nza, int hacksz, int noffs,
fprintf(stderr,"From routine : %s : %d \n","writeHllDeviceDoubleComplex",i);
}
return SPGPU_SUCCESS;
#else
return SPGPU_UNSUPPORTED;
#endif
}
#endif

5
cuda/hlldev.h
Unescape Escape View File

@ -32,7 +32,6 @@
#ifndef _HLLDEV_H_
#define _HLLDEV_H_
#ifdef HAVE_SPGPU
#include "cintrf.h"
#include "hell.h"
@ -154,8 +153,4 @@ int psi_cuda_z_CopyCooToHlg(spgpuHandle_t handle,int nr, int nc, int nza,
int *idiag, int *rP, double complex *cM);
#else
#define CINTRF_UNSUPPORTED -1
#endif
#endif

5
cuda/hlldev_mod.F90
Unescape Escape View File

@ -43,8 +43,6 @@ module hlldev_mod
integer(c_int) :: firstIndex
end type hlldev_parms
#ifdef HAVE_SPGPU
interface
function bldHllDeviceParams(hksize, rows, nzeros, allocsize, elementType, firstIndex) &
& result(res) bind(c,name='bldHllDeviceParams')
@ -267,7 +265,4 @@ module hlldev_mod
end interface
#endif
end module hlldev_mod

6
cuda/impl/psb_c_cuda_cp_csrg_from_coo.F90
Unescape Escape View File

@ -32,12 +32,8 @@
subroutine psb_c_cuda_cp_csrg_from_coo(a,b,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use cusparse_mod
use psb_c_cuda_csrg_mat_mod, psb_protect_name => psb_c_cuda_cp_csrg_from_coo
#else
use psb_c_cuda_csrg_mat_mod
#endif
implicit none
class(psb_c_cuda_csrg_sparse_mat), intent(inout) :: a
@ -48,10 +44,8 @@ subroutine psb_c_cuda_cp_csrg_from_coo(a,b,info)
call a%psb_c_csr_sparse_mat%cp_from_coo(b,info)
if (info /= 0) goto 9999
#ifdef HAVE_SPGPU
call a%to_gpu(info)
if (info /= 0) goto 9999
#endif
return

6
cuda/impl/psb_c_cuda_cp_csrg_from_fmt.F90
Unescape Escape View File

@ -32,12 +32,8 @@
subroutine psb_c_cuda_cp_csrg_from_fmt(a,b,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use cusparse_mod
use psb_c_cuda_csrg_mat_mod, psb_protect_name => psb_c_cuda_cp_csrg_from_fmt
#else
use psb_c_cuda_csrg_mat_mod
#endif
!use iso_c_binding
implicit none
@ -53,9 +49,7 @@ subroutine psb_c_cuda_cp_csrg_from_fmt(a,b,info)
class default
call a%psb_c_csr_sparse_mat%cp_from_fmt(b,info)
if (info /= 0) return
#ifdef HAVE_SPGPU
call a%to_gpu(info)
#endif
end select
end subroutine psb_c_cuda_cp_csrg_from_fmt

6
cuda/impl/psb_c_cuda_cp_diag_from_coo.F90
Unescape Escape View File

@ -33,13 +33,9 @@
subroutine psb_c_cuda_cp_diag_from_coo(a,b,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use diagdev_mod
use psb_vectordev_mod
use psb_c_cuda_diag_mat_mod, psb_protect_name => psb_c_cuda_cp_diag_from_coo
#else
use psb_c_cuda_diag_mat_mod
#endif
implicit none
class(psb_c_cuda_diag_sparse_mat), intent(inout) :: a
@ -50,10 +46,8 @@ subroutine psb_c_cuda_cp_diag_from_coo(a,b,info)
info = psb_success_
call a%psb_c_dia_sparse_mat%cp_from_coo(b,info)
#ifdef HAVE_SPGPU
call a%to_gpu(info)
if (info /= 0) goto 9999
#endif
return

23
cuda/impl/psb_c_cuda_cp_elg_from_coo.F90
Unescape Escape View File

@ -28,20 +28,14 @@
! ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_cp_elg_from_coo(a,b,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use elldev_mod
use psb_vectordev_mod
use psb_c_cuda_elg_mat_mod, psb_protect_name => psb_c_cuda_cp_elg_from_coo
use psi_ext_util_mod
use psb_cuda_env_mod
#else
use psb_c_cuda_elg_mat_mod
#endif
implicit none
class(psb_c_cuda_elg_sparse_mat), intent(inout) :: a
@ -57,16 +51,11 @@ subroutine psb_c_cuda_cp_elg_from_coo(a,b,info)
integer(psb_ipk_), allocatable :: idisp(:)
info = psb_success_
#ifdef HAVE_SPGPU
hacksize = max(1,psb_cuda_WarpSize())
#else
hacksize = 1
#endif
if (b%is_dev()) call b%sync()
if (b%is_by_rows()) then
#ifdef HAVE_SPGPU
call psi_c_count_ell_from_coo(a,b,idisp,ldv,nzm,info,hacksize=hacksize)
@ -82,15 +71,8 @@ subroutine psb_c_cuda_cp_elg_from_coo(a,b,info)
if (info == 0) info = psi_CopyCooToElg(nr,nc,nza, hacksize,ldv,nzm, &
& a%irn,idisp,b%ja,b%val, a%deviceMat)
call a%set_dev()
#else
call psi_c_convert_ell_from_coo(a,b,info,hacksize=hacksize)
call a%set_host()
#endif
else
call b%cp_to_coo(tmp,info)
#ifdef HAVE_SPGPU
call psi_c_count_ell_from_coo(a,tmp,idisp,ldv,nzm,info,hacksize=hacksize)
@ -107,11 +89,6 @@ subroutine psb_c_cuda_cp_elg_from_coo(a,b,info)
& a%irn,idisp,tmp%ja,tmp%val, a%deviceMat)
call a%set_dev()
#else
call psi_c_convert_ell_from_coo(a,tmp,info,hacksize=hacksize)
call a%set_host()
#endif
end if
if (info /= psb_success_) goto 9999

12
cuda/impl/psb_c_cuda_cp_elg_from_fmt.F90
Unescape Escape View File

@ -33,13 +33,9 @@
subroutine psb_c_cuda_cp_elg_from_fmt(a,b,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use elldev_mod
use psb_vectordev_mod
use psb_c_cuda_elg_mat_mod, psb_protect_name => psb_c_cuda_cp_elg_from_fmt
#else
use psb_c_cuda_elg_mat_mod
#endif
implicit none
class(psb_c_cuda_elg_sparse_mat), intent(inout) :: a
@ -51,9 +47,7 @@ subroutine psb_c_cuda_cp_elg_from_fmt(a,b,info)
Integer(Psb_ipk_) :: nza, nr, i,j,irw, idl,err_act, nc, ld, nzm, m
integer(psb_ipk_) :: debug_level, debug_unit
character(len=20) :: name
#ifdef HAVE_SPGPU
type(elldev_parms) :: gpu_parms
#endif
info = psb_success_
if (b%is_dev()) call b%sync()
@ -67,13 +61,9 @@ subroutine psb_c_cuda_cp_elg_from_fmt(a,b,info)
m = b%get_nrows()
nc = b%get_ncols()
nza = b%get_nzeros()
#ifdef HAVE_SPGPU
gpu_parms = FgetEllDeviceParams(m,nzm,nza,nc,spgpu_type_double,1)
ld = gpu_parms%pitch
nzm = gpu_parms%maxRowSize
#else
ld = m
#endif
a%psb_c_base_sparse_mat = b%psb_c_base_sparse_mat
if (info == 0) call psb_safe_cpy( b%idiag, a%idiag , info)
if (info == 0) call psb_safe_cpy( b%irn, a%irn , info)
@ -88,9 +78,7 @@ subroutine psb_c_cuda_cp_elg_from_fmt(a,b,info)
a%val(1:m,1:nzm) = b%val(1:m,1:nzm)
end if
a%nzt = nza
#ifdef HAVE_SPGPU
call a%to_gpu(info)
#endif
class default

10
cuda/impl/psb_c_cuda_cp_hdiag_from_coo.F90
Unescape Escape View File

@ -28,19 +28,13 @@
! ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_cp_hdiag_from_coo(a,b,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use hdiagdev_mod
use psb_vectordev_mod
use psb_c_cuda_hdiag_mat_mod, psb_protect_name => psb_c_cuda_cp_hdiag_from_coo
use psb_cuda_env_mod
#else
use psb_c_cuda_hdiag_mat_mod
#endif
implicit none
class(psb_c_cuda_hdiag_sparse_mat), intent(inout) :: a
@ -53,16 +47,12 @@ subroutine psb_c_cuda_cp_hdiag_from_coo(a,b,info)
info = psb_success_
#ifdef HAVE_SPGPU
a%hacksize = psb_cuda_WarpSize()
#endif
call a%psb_c_hdia_sparse_mat%cp_from_coo(b,info)
#ifdef HAVE_SPGPU
call a%to_gpu(info)
if (info /= 0) goto 9999
#endif
return

8
cuda/impl/psb_c_cuda_cp_hlg_from_coo.F90
Unescape Escape View File

@ -33,14 +33,10 @@
subroutine psb_c_cuda_cp_hlg_from_coo(a,b,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use hlldev_mod
use psb_vectordev_mod
use psb_cuda_env_mod
use psb_c_cuda_hlg_mat_mod, psb_protect_name => psb_c_cuda_cp_hlg_from_coo
#else
use psb_c_cuda_hlg_mat_mod
#endif
implicit none
class(psb_c_cuda_hlg_sparse_mat), intent(inout) :: a
@ -61,11 +57,7 @@ subroutine psb_c_cuda_cp_hlg_from_coo(a,b,info)
info = psb_success_
debug_unit = psb_get_debug_unit()
debug_level = psb_get_debug_level()
#ifdef HAVE_SPGPU
hksz = max(1,psb_cuda_WarpSize())
#else
hksz = psi_get_hksz()
#endif
if (b%is_by_rows()) then

6
cuda/impl/psb_c_cuda_cp_hlg_from_fmt.F90
Unescape Escape View File

@ -33,13 +33,9 @@
subroutine psb_c_cuda_cp_hlg_from_fmt(a,b,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use hlldev_mod
use psb_vectordev_mod
use psb_c_cuda_hlg_mat_mod, psb_protect_name => psb_c_cuda_cp_hlg_from_fmt
#else
use psb_c_cuda_hlg_mat_mod
#endif
implicit none
class(psb_c_cuda_hlg_sparse_mat), intent(inout) :: a
@ -53,9 +49,7 @@ subroutine psb_c_cuda_cp_hlg_from_fmt(a,b,info)
call a%cp_from_coo(b,info)
class default
call a%psb_c_hll_sparse_mat%cp_from_fmt(b,info)
#ifdef HAVE_SPGPU
if (info == 0) call a%to_gpu(info)
#endif
end select
if (info /= 0) goto 9999

6
cuda/impl/psb_c_cuda_cp_hybg_from_coo.F90
Unescape Escape View File

@ -33,12 +33,8 @@
subroutine psb_c_cuda_cp_hybg_from_coo(a,b,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use cusparse_mod
use psb_c_cuda_hybg_mat_mod, psb_protect_name => psb_c_cuda_cp_hybg_from_coo
#else
use psb_c_cuda_hybg_mat_mod
#endif
implicit none
class(psb_c_cuda_hybg_sparse_mat), intent(inout) :: a
@ -49,10 +45,8 @@ subroutine psb_c_cuda_cp_hybg_from_coo(a,b,info)
call a%psb_c_csr_sparse_mat%cp_from_coo(b,info)
if (info /= 0) goto 9999
#ifdef HAVE_SPGPU
call a%to_gpu(info)
if (info /= 0) goto 9999
#endif
return

6
cuda/impl/psb_c_cuda_cp_hybg_from_fmt.F90
Unescape Escape View File

@ -33,12 +33,8 @@
subroutine psb_c_cuda_cp_hybg_from_fmt(a,b,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use cusparse_mod
use psb_c_cuda_hybg_mat_mod, psb_protect_name => psb_c_cuda_cp_hybg_from_fmt
#else
use psb_c_cuda_hybg_mat_mod
#endif
implicit none
class(psb_c_cuda_hybg_sparse_mat), intent(inout) :: a
@ -53,9 +49,7 @@ subroutine psb_c_cuda_cp_hybg_from_fmt(a,b,info)
class default
call a%psb_c_csr_sparse_mat%cp_from_fmt(b,info)
if (info /= 0) return
#ifdef HAVE_SPGPU
call a%to_gpu(info)
#endif
end select
end subroutine psb_c_cuda_cp_hybg_from_fmt

6
cuda/impl/psb_c_cuda_csrg_allocate_mnnz.F90
Unescape Escape View File

@ -33,12 +33,8 @@
subroutine psb_c_cuda_csrg_allocate_mnnz(m,n,a,nz)
use psb_base_mod
#ifdef HAVE_SPGPU
use cusparse_mod
use psb_c_cuda_csrg_mat_mod, psb_protect_name => psb_c_cuda_csrg_allocate_mnnz
#else
use psb_c_cuda_csrg_mat_mod
#endif
implicit none
integer(psb_ipk_), intent(in) :: m,n
class(psb_c_cuda_csrg_sparse_mat), intent(inout) :: a
@ -52,11 +48,9 @@ subroutine psb_c_cuda_csrg_allocate_mnnz(m,n,a,nz)
call a%psb_c_csr_sparse_mat%allocate(m,n,nz)
#ifdef HAVE_SPGPU
info = initFcusparse()
if (info == 0) call a%to_gpu(info,nzrm=nz)
if (info /= 0) goto 9999
#endif
call psb_erractionrestore(err_act)
return

8
cuda/impl/psb_c_cuda_csrg_csmm.F90
Unescape Escape View File

@ -33,14 +33,10 @@
subroutine psb_c_cuda_csrg_csmm(alpha,a,x,beta,y,info,trans)
use psb_base_mod
#ifdef HAVE_SPGPU
use cusparse_mod
use elldev_mod
use psb_vectordev_mod
use psb_c_cuda_csrg_mat_mod, psb_protect_name => psb_c_cuda_csrg_csmm
#else
use psb_c_cuda_csrg_mat_mod
#endif
implicit none
class(psb_c_cuda_csrg_sparse_mat), intent(in) :: a
complex(psb_spk_), intent(in) :: alpha, beta, x(:,:)
@ -94,7 +90,6 @@ subroutine psb_c_cuda_csrg_csmm(alpha,a,x,beta,y,info,trans)
end if
#ifdef HAVE_SPGPU
if (tra) then
call a%psb_c_csr_sparse_mat%spmm(alpha,x,beta,y,info,trans)
else
@ -119,9 +114,6 @@ subroutine psb_c_cuda_csrg_csmm(alpha,a,x,beta,y,info,trans)
call freeMultiVecDevice(gpX)
call freeMultiVecDevice(gpY)
endif
#else
call a%psb_c_csr_sparse_mat%spmm(alpha,x,beta,y,info,trans)
#endif
call psb_erractionrestore(err_act)
return

8
cuda/impl/psb_c_cuda_csrg_csmv.F90
Unescape Escape View File

@ -33,14 +33,10 @@
subroutine psb_c_cuda_csrg_csmv(alpha,a,x,beta,y,info,trans)
use psb_base_mod
#ifdef HAVE_SPGPU
use cusparse_mod
use elldev_mod
use psb_vectordev_mod
use psb_c_cuda_csrg_mat_mod, psb_protect_name => psb_c_cuda_csrg_csmv
#else
use psb_c_cuda_csrg_mat_mod
#endif
implicit none
class(psb_c_cuda_csrg_sparse_mat), intent(in) :: a
complex(psb_spk_), intent(in) :: alpha, beta, x(:)
@ -96,7 +92,6 @@ subroutine psb_c_cuda_csrg_csmv(alpha,a,x,beta,y,info,trans)
goto 9999
end if
#ifdef HAVE_SPGPU
if (tra) then
call a%psb_c_csr_sparse_mat%spmm(alpha,x,beta,y,info,trans)
else
@ -124,9 +119,6 @@ subroutine psb_c_cuda_csrg_csmv(alpha,a,x,beta,y,info,trans)
call freeMultiVecDevice(gpX)
call freeMultiVecDevice(gpY)
endif
#else
call a%psb_c_csr_sparse_mat%spmm(alpha,x,beta,y,info,trans)
#endif
call psb_erractionrestore(err_act)
return

6
cuda/impl/psb_c_cuda_csrg_from_gpu.F90
Unescape Escape View File

@ -33,13 +33,9 @@
subroutine psb_c_cuda_csrg_from_gpu(a,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use elldev_mod
use psb_vectordev_mod
use psb_c_cuda_csrg_mat_mod, psb_protect_name => psb_c_cuda_csrg_from_gpu
#else
use psb_c_cuda_csrg_mat_mod
#endif
implicit none
class(psb_c_cuda_csrg_sparse_mat), intent(inout) :: a
integer(psb_ipk_), intent(out) :: info
@ -48,7 +44,6 @@ subroutine psb_c_cuda_csrg_from_gpu(a,info)
info = 0
#ifdef HAVE_SPGPU
if (.not.(c_associated(a%deviceMat%mat))) then
call a%free()
return
@ -68,6 +63,5 @@ subroutine psb_c_cuda_csrg_from_gpu(a,info)
#endif
call a%set_sync()
#endif
end subroutine psb_c_cuda_csrg_from_gpu

11
cuda/impl/psb_c_cuda_csrg_inner_vect_sv.F90
Unescape Escape View File

@ -32,13 +32,9 @@
subroutine psb_c_cuda_csrg_inner_vect_sv(alpha,a,x,beta,y,info,trans)
use psb_base_mod
#ifdef HAVE_SPGPU
use elldev_mod
use psb_vectordev_mod
use psb_c_cuda_csrg_mat_mod, psb_protect_name => psb_c_cuda_csrg_inner_vect_sv
#else
use psb_c_cuda_csrg_mat_mod
#endif
use psb_c_cuda_vect_mod
implicit none
class(psb_c_cuda_csrg_sparse_mat), intent(in) :: a
@ -75,7 +71,6 @@ subroutine psb_c_cuda_csrg_inner_vect_sv(alpha,a,x,beta,y,info,trans)
tra = (psb_toupper(trans_) == 'T').or.(psb_toupper(trans_)=='C')
#ifdef HAVE_SPGPU
if (tra.or.(beta/=dzero)) then
call x%sync()
call y%sync()
@ -112,12 +107,6 @@ subroutine psb_c_cuda_csrg_inner_vect_sv(alpha,a,x,beta,y,info,trans)
call y%bld(ry)
end select
end if
#else
call x%sync()
call y%sync()
call a%psb_c_csr_sparse_mat%inner_spsm(alpha,x,beta,y,info,trans)
call y%set_host()
#endif
if (info /= psb_success_) then
info = psb_err_from_subroutine_
call psb_errpush(info,name, a_err='csrg_vect_sv')

6
cuda/impl/psb_c_cuda_csrg_reallocate_nz.F90
Unescape Escape View File

@ -33,12 +33,8 @@
subroutine psb_c_cuda_csrg_reallocate_nz(nz,a)
use psb_base_mod
#ifdef HAVE_SPGPU
use cusparse_mod
use psb_c_cuda_csrg_mat_mod, psb_protect_name => psb_c_cuda_csrg_reallocate_nz
#else
use psb_c_cuda_csrg_mat_mod
#endif
implicit none
integer(psb_ipk_), intent(in) :: nz
class(psb_c_cuda_csrg_sparse_mat), intent(inout) :: a
@ -55,10 +51,8 @@ subroutine psb_c_cuda_csrg_reallocate_nz(nz,a)
!
call a%psb_c_csr_sparse_mat%reallocate(nz)
#ifdef HAVE_SPGPU
call a%to_gpu(info,nzrm=nz)
if (info /= 0) goto 9999
#endif
call psb_erractionrestore(err_act)
return

6
cuda/impl/psb_c_cuda_csrg_scal.F90
Unescape Escape View File

@ -33,12 +33,8 @@
subroutine psb_c_cuda_csrg_scal(d,a,info,side)
use psb_base_mod
#ifdef HAVE_SPGPU
use cusparse_mod
use psb_c_cuda_csrg_mat_mod, psb_protect_name => psb_c_cuda_csrg_scal
#else
use psb_c_cuda_csrg_mat_mod
#endif
implicit none
class(psb_c_cuda_csrg_sparse_mat), intent(inout) :: a
complex(psb_spk_), intent(in) :: d(:)
@ -58,10 +54,8 @@ subroutine psb_c_cuda_csrg_scal(d,a,info,side)
call a%psb_c_csr_sparse_mat%scal(d,info,side=side)
if (info /= 0) goto 9999
#ifdef HAVE_SPGPU
call a%to_gpu(info)
if (info /= 0) goto 9999
#endif
call psb_erractionrestore(err_act)
return

6
cuda/impl/psb_c_cuda_csrg_scals.F90
Unescape Escape View File

@ -33,12 +33,8 @@
subroutine psb_c_cuda_csrg_scals(d,a,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use cusparse_mod
use psb_c_cuda_csrg_mat_mod, psb_protect_name => psb_c_cuda_csrg_scals
#else
use psb_c_cuda_csrg_mat_mod
#endif
implicit none
class(psb_c_cuda_csrg_sparse_mat), intent(inout) :: a
complex(psb_spk_), intent(in) :: d
@ -56,10 +52,8 @@ subroutine psb_c_cuda_csrg_scals(d,a,info)
if (info /= 0) goto 9999
#ifdef HAVE_SPGPU
call a%to_gpu(info)
if (info /= 0) goto 9999
#endif
call psb_erractionrestore(err_act)
return

6
cuda/impl/psb_c_cuda_csrg_to_gpu.F90
Unescape Escape View File

@ -33,12 +33,8 @@
subroutine psb_c_cuda_csrg_to_gpu(a,info,nzrm)
use psb_base_mod
#ifdef HAVE_SPGPU
use cusparse_mod
use psb_c_cuda_csrg_mat_mod, psb_protect_name => psb_c_cuda_csrg_to_gpu
#else
use psb_c_cuda_csrg_mat_mod
#endif
implicit none
class(psb_c_cuda_csrg_sparse_mat), intent(inout) :: a
integer(psb_ipk_), intent(out) :: info
@ -51,7 +47,6 @@ subroutine psb_c_cuda_csrg_to_gpu(a,info,nzrm)
info = 0
#ifdef HAVE_SPGPU
if ((.not.allocated(a%val)).or.(.not.allocated(a%ja))) return
m = a%get_nrows()
@ -320,6 +315,5 @@ subroutine psb_c_cuda_csrg_to_gpu(a,info,nzrm)
if (info /= 0) then
write(0,*) 'Error in CSRG_TO_GPU ',info
end if
#endif
end subroutine psb_c_cuda_csrg_to_gpu

8
cuda/impl/psb_c_cuda_csrg_vect_mv.F90
Unescape Escape View File

@ -33,14 +33,10 @@
subroutine psb_c_cuda_csrg_vect_mv(alpha,a,x,beta,y,info,trans)
use psb_base_mod
#ifdef HAVE_SPGPU
use cusparse_mod
use elldev_mod
use psb_vectordev_mod
use psb_c_cuda_csrg_mat_mod, psb_protect_name => psb_c_cuda_csrg_vect_mv
#else
use psb_c_cuda_csrg_mat_mod
#endif
use psb_c_cuda_vect_mod
implicit none
class(psb_c_cuda_csrg_sparse_mat), intent(in) :: a
@ -72,7 +68,6 @@ subroutine psb_c_cuda_csrg_vect_mv(alpha,a,x,beta,y,info,trans)
tra = (psb_toupper(trans_) == 'T').or.(psb_toupper(trans_)=='C')
#ifdef HAVE_SPGPU
if (tra) then
if (.not.x%is_host()) call x%sync()
if (beta /= czero) then
@ -112,9 +107,6 @@ subroutine psb_c_cuda_csrg_vect_mv(alpha,a,x,beta,y,info,trans)
call y%bld(ry)
end select
end if
#else
call a%psb_c_csr_sparse_mat%spmm(alpha,x,beta,y,info,trans)
#endif
if (info /= 0) goto 9999
call psb_erractionrestore(err_act)
return

9
cuda/impl/psb_c_cuda_diag_csmv.F90
Unescape Escape View File

@ -33,13 +33,9 @@
subroutine psb_c_cuda_diag_csmv(alpha,a,x,beta,y,info,trans)
use psb_base_mod
#ifdef HAVE_SPGPU
use diagdev_mod
use psb_vectordev_mod
use psb_c_cuda_diag_mat_mod, psb_protect_name => psb_c_cuda_diag_csmv
#else
use psb_c_cuda_diag_mat_mod
#endif
implicit none
class(psb_c_cuda_diag_sparse_mat), intent(in) :: a
complex(psb_spk_), intent(in) :: alpha, beta, x(:)
@ -94,7 +90,6 @@ subroutine psb_c_cuda_diag_csmv(alpha,a,x,beta,y,info,trans)
goto 9999
end if
#ifdef HAVE_SPGPU
if (tra) then
call a%psb_c_dia_sparse_mat%spmm(alpha,x,beta,y,info,trans)
else
@ -121,9 +116,6 @@ subroutine psb_c_cuda_diag_csmv(alpha,a,x,beta,y,info,trans)
call freeMultiVecDevice(gpX)
call freeMultiVecDevice(gpY)
endif
#else
call a%psb_c_dia_sparse_mat%spmm(alpha,x,beta,y,info,trans)
#endif
call psb_erractionrestore(err_act)
return
@ -132,5 +124,4 @@ subroutine psb_c_cuda_diag_csmv(alpha,a,x,beta,y,info,trans)
return
end subroutine psb_c_cuda_diag_csmv

8
cuda/impl/psb_c_cuda_diag_to_gpu.F90
Unescape Escape View File

@ -33,13 +33,9 @@
subroutine psb_c_cuda_diag_to_gpu(a,info,nzrm)
use psb_base_mod
#ifdef HAVE_SPGPU
use diagdev_mod
use psb_vectordev_mod
use psb_c_cuda_diag_mat_mod, psb_protect_name => psb_c_cuda_diag_to_gpu
#else
use psb_c_cuda_diag_mat_mod
#endif
use iso_c_binding
implicit none
class(psb_c_cuda_diag_sparse_mat), intent(inout) :: a
@ -47,13 +43,10 @@ subroutine psb_c_cuda_diag_to_gpu(a,info,nzrm)
integer(psb_ipk_), intent(in), optional :: nzrm
integer(psb_ipk_) :: m, nzm, n, c,pitch,maxrowsize,d
#ifdef HAVE_SPGPU
type(diagdev_parms) :: gpu_parms
#endif
info = 0
#ifdef HAVE_SPGPU
if ((.not.allocated(a%data)).or.(.not.allocated(a%offset))) return
n = size(a%data,1)
@ -69,6 +62,5 @@ subroutine psb_c_cuda_diag_to_gpu(a,info,nzrm)
if (info == 0) info = &
& writeDiagDevice(a%deviceMat,a%data,a%offset,n)
! if (info /= 0) goto 9999
#endif
end subroutine psb_c_cuda_diag_to_gpu

10
cuda/impl/psb_c_cuda_diag_vect_mv.F90
Unescape Escape View File

@ -28,18 +28,12 @@
! ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_diag_vect_mv(alpha,a,x,beta,y,info,trans)
use psb_base_mod
#ifdef HAVE_SPGPU
use diagdev_mod
use psb_vectordev_mod
use psb_c_cuda_diag_mat_mod, psb_protect_name => psb_c_cuda_diag_vect_mv
#else
use psb_c_cuda_diag_mat_mod
#endif
use psb_c_cuda_vect_mod
implicit none
class(psb_c_cuda_diag_sparse_mat), intent(in) :: a
@ -71,7 +65,6 @@ subroutine psb_c_cuda_diag_vect_mv(alpha,a,x,beta,y,info,trans)
tra = (psb_toupper(trans_) == 'T').or.(psb_toupper(trans_)=='C')
#ifdef HAVE_SPGPU
if (tra) then
if (.not.x%is_host()) call x%sync()
if (beta /= szero) then
@ -112,9 +105,6 @@ subroutine psb_c_cuda_diag_vect_mv(alpha,a,x,beta,y,info,trans)
end select
end if
#else
call a%psb_c_dia_sparse_mat%spmm(alpha,x,beta,y,info,trans)
#endif
if (info /= 0) goto 9999
call psb_erractionrestore(err_act)
return

33
cuda/impl/psb_c_cuda_dnsg_mat_impl.F90
Unescape Escape View File

@ -32,13 +32,9 @@
subroutine psb_c_cuda_dnsg_vect_mv(alpha,a,x,beta,y,info,trans)
use psb_base_mod
use psb_c_cuda_vect_mod
#ifdef HAVE_SPGPU
use dnsdev_mod
use psb_c_vectordev_mod
use psb_c_cuda_dnsg_mat_mod, psb_protect_name => psb_c_cuda_dnsg_vect_mv
#else
use psb_c_cuda_dnsg_mat_mod
#endif
implicit none
class(psb_c_cuda_dnsg_sparse_mat), intent(in) :: a
complex(psb_spk_), intent(in) :: alpha, beta
@ -123,13 +119,9 @@ end subroutine psb_c_cuda_dnsg_vect_mv
subroutine psb_c_cuda_dnsg_mold(a,b,info)
use psb_base_mod
use psb_c_cuda_vect_mod
#ifdef HAVE_SPGPU
use dnsdev_mod
use psb_c_vectordev_mod
use psb_c_cuda_dnsg_mat_mod, psb_protect_name => psb_c_cuda_dnsg_mold
#else
use psb_c_cuda_dnsg_mat_mod
#endif
implicit none
class(psb_c_cuda_dnsg_sparse_mat), intent(in) :: a
class(psb_c_base_sparse_mat), intent(inout), allocatable :: b
@ -190,17 +182,12 @@ end subroutine psb_c_cuda_dnsg_mold
!!$ end subroutine psb_c_cuda_dnsg_allocate_mnnz
!!$ end interface
subroutine psb_c_cuda_dnsg_to_gpu(a,info)
use psb_base_mod
use psb_c_cuda_vect_mod
#ifdef HAVE_SPGPU
use dnsdev_mod
use psb_c_vectordev_mod
use psb_c_cuda_dnsg_mat_mod, psb_protect_name => psb_c_cuda_dnsg_to_gpu
#else
use psb_c_cuda_dnsg_mat_mod
#endif
class(psb_c_cuda_dnsg_sparse_mat), intent(inout) :: a
integer(psb_ipk_), intent(out) :: info
Integer(Psb_ipk_) :: err_act, pitch, lda
@ -209,15 +196,12 @@ subroutine psb_c_cuda_dnsg_to_gpu(a,info)
call psb_erractionsave(err_act)
info = psb_success_
#ifdef HAVE_SPGPU
if (debug) write(0,*) 'DNS_TO_GPU',size(a%val,1),size(a%val,2)
info = FallocDnsDevice(a%deviceMat,a%get_nrows(),a%get_ncols(),&
& spgpu_type_complex_float,1)
if (info == 0) info = writeDnsDevice(a%deviceMat,a%val,size(a%val,1),size(a%val,2))
if (debug) write(0,*) 'DNS_TO_GPU: From writeDnsDEvice',info
#endif
if (info /= 0) goto 9999
call psb_erractionrestore(err_act)
return
@ -233,13 +217,9 @@ end subroutine psb_c_cuda_dnsg_to_gpu
subroutine psb_c_cuda_cp_dnsg_from_coo(a,b,info)
use psb_base_mod
use psb_c_cuda_vect_mod
#ifdef HAVE_SPGPU
use dnsdev_mod
use psb_c_vectordev_mod
use psb_c_cuda_dnsg_mat_mod, psb_protect_name => psb_c_cuda_cp_dnsg_from_coo
#else
use psb_c_cuda_dnsg_mat_mod
#endif
implicit none
class(psb_c_cuda_dnsg_sparse_mat), intent(inout) :: a
@ -272,13 +252,9 @@ end subroutine psb_c_cuda_cp_dnsg_from_coo
subroutine psb_c_cuda_cp_dnsg_from_fmt(a,b,info)
use psb_base_mod
use psb_c_cuda_vect_mod
#ifdef HAVE_SPGPU
use dnsdev_mod
use psb_c_vectordev_mod
use psb_c_cuda_dnsg_mat_mod, psb_protect_name => psb_c_cuda_cp_dnsg_from_fmt
#else
use psb_c_cuda_dnsg_mat_mod
#endif
implicit none
class(psb_c_cuda_dnsg_sparse_mat), intent(inout) :: a
@ -348,13 +324,9 @@ end subroutine psb_c_cuda_cp_dnsg_from_fmt
subroutine psb_c_cuda_mv_dnsg_from_coo(a,b,info)
use psb_base_mod
use psb_c_cuda_vect_mod
#ifdef HAVE_SPGPU
use dnsdev_mod
use psb_c_vectordev_mod
use psb_c_cuda_dnsg_mat_mod, psb_protect_name => psb_c_cuda_mv_dnsg_from_coo
#else
use psb_c_cuda_dnsg_mat_mod
#endif
implicit none
class(psb_c_cuda_dnsg_sparse_mat), intent(inout) :: a
@ -384,17 +356,12 @@ subroutine psb_c_cuda_mv_dnsg_from_coo(a,b,info)
end subroutine psb_c_cuda_mv_dnsg_from_coo
subroutine psb_c_cuda_mv_dnsg_from_fmt(a,b,info)
use psb_base_mod
use psb_c_cuda_vect_mod
#ifdef HAVE_SPGPU
use dnsdev_mod
use psb_c_vectordev_mod
use psb_c_cuda_dnsg_mat_mod, psb_protect_name => psb_c_cuda_mv_dnsg_from_fmt
#else
use psb_c_cuda_dnsg_mat_mod
#endif
implicit none
class(psb_c_cuda_dnsg_sparse_mat), intent(inout) :: a
class(psb_c_base_sparse_mat), intent(inout) :: b

14
cuda/impl/psb_c_cuda_elg_allocate_mnnz.F90
Unescape Escape View File

@ -28,18 +28,12 @@
! ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_elg_allocate_mnnz(m,n,a,nz)
use psb_base_mod
#ifdef HAVE_SPGPU
use elldev_mod
use psb_vectordev_mod
use psb_c_cuda_elg_mat_mod, psb_protect_name => psb_c_cuda_elg_allocate_mnnz
#else
use psb_c_cuda_elg_mat_mod
#endif
implicit none
integer(psb_ipk_), intent(in) :: m,n
class(psb_c_cuda_elg_sparse_mat), intent(inout) :: a
@ -47,9 +41,7 @@ subroutine psb_c_cuda_elg_allocate_mnnz(m,n,a,nz)
Integer(Psb_ipk_) :: err_act, info, nz_,ld
character(len=20) :: name='allocate_mnz'
logical, parameter :: debug=.false.
#ifdef HAVE_SPGPU
type(elldev_parms) :: gpu_parms
#endif
call psb_erractionsave(err_act)
info = psb_success_
@ -74,13 +66,9 @@ subroutine psb_c_cuda_elg_allocate_mnnz(m,n,a,nz)
goto 9999
endif
#ifdef HAVE_SPGPU
gpu_parms = FgetEllDeviceParams(m,nz_,nz_*m,n,spgpu_type_complex_float,1)
ld = gpu_parms%pitch
nz_ = gpu_parms%maxRowSize
#else
ld = m
#endif
if (info == psb_success_) call psb_realloc(m,a%irn,info)
if (info == psb_success_) call psb_realloc(m,a%idiag,info)
@ -98,10 +86,8 @@ subroutine psb_c_cuda_elg_allocate_mnnz(m,n,a,nz)
call a%set_dupl(psb_dupl_def_)
end if
#ifdef HAVE_SPGPU
call a%to_gpu(info,nzrm=nz_)
if (info /= 0) goto 9999
#endif
call psb_erractionrestore(err_act)
return

1
cuda/impl/psb_c_cuda_elg_asb.f90
Unescape Escape View File

@ -29,7 +29,6 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_elg_asb(a)
use psb_base_mod

10
cuda/impl/psb_c_cuda_elg_csmm.F90
Unescape Escape View File

@ -29,17 +29,12 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_elg_csmm(alpha,a,x,beta,y,info,trans)
use psb_base_mod
#ifdef HAVE_SPGPU
use elldev_mod
use psb_vectordev_mod
use psb_c_cuda_elg_mat_mod, psb_protect_name => psb_c_cuda_elg_csmm
#else
use psb_c_cuda_elg_mat_mod
#endif
implicit none
class(psb_c_cuda_elg_sparse_mat), intent(in) :: a
complex(psb_spk_), intent(in) :: alpha, beta, x(:,:)
@ -92,8 +87,6 @@ subroutine psb_c_cuda_elg_csmm(alpha,a,x,beta,y,info,trans)
goto 9999
end if
#ifdef HAVE_SPGPU
if (tra) then
if (a%is_dev()) call a%sync()
call a%psb_c_ell_sparse_mat%spmm(alpha,x,beta,y,info,trans)
@ -119,9 +112,6 @@ subroutine psb_c_cuda_elg_csmm(alpha,a,x,beta,y,info,trans)
call freeMultiVecDevice(gpX)
call freeMultiVecDevice(gpY)
endif
#else
call a%psb_c_ell_sparse_mat%spmm(alpha,x,beta,y,info,trans)
#endif
call psb_erractionrestore(err_act)
return

9
cuda/impl/psb_c_cuda_elg_csmv.F90
Unescape Escape View File

@ -29,17 +29,12 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_elg_csmv(alpha,a,x,beta,y,info,trans)
use psb_base_mod
#ifdef HAVE_SPGPU
use elldev_mod
use psb_vectordev_mod
use psb_c_cuda_elg_mat_mod, psb_protect_name => psb_c_cuda_elg_csmv
#else
use psb_c_cuda_elg_mat_mod
#endif
implicit none
class(psb_c_cuda_elg_sparse_mat), intent(in) :: a
complex(psb_spk_), intent(in) :: alpha, beta, x(:)
@ -94,7 +89,6 @@ subroutine psb_c_cuda_elg_csmv(alpha,a,x,beta,y,info,trans)
goto 9999
end if
#ifdef HAVE_SPGPU
if (tra) then
if (a%is_dev()) call a%sync()
call a%psb_c_ell_sparse_mat%spmm(alpha,x,beta,y,info,trans)
@ -122,9 +116,6 @@ subroutine psb_c_cuda_elg_csmv(alpha,a,x,beta,y,info,trans)
call freeMultiVecDevice(gpX)
call freeMultiVecDevice(gpY)
endif
#else
call a%psb_c_ell_sparse_mat%spmm(alpha,x,beta,y,info,trans)
#endif
call psb_erractionrestore(err_act)
return

9
cuda/impl/psb_c_cuda_elg_csput.F90
Unescape Escape View File

@ -29,17 +29,12 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_elg_csput_a(nz,ia,ja,val,a,imin,imax,jmin,jmax,info)
use psb_base_mod
use iso_c_binding
#ifdef HAVE_SPGPU
use elldev_mod
use psb_c_cuda_elg_mat_mod, psb_protect_name => psb_c_cuda_elg_csput_a
#else
use psb_c_cuda_elg_mat_mod
#endif
implicit none
class(psb_c_cuda_elg_sparse_mat), intent(inout) :: a
@ -128,13 +123,9 @@ subroutine psb_c_cuda_elg_csput_v(nz,ia,ja,val,a,imin,imax,jmin,jmax,info)
use psb_base_mod
use iso_c_binding
#ifdef HAVE_SPGPU
use elldev_mod
use psb_c_cuda_elg_mat_mod, psb_protect_name => psb_c_cuda_elg_csput_v
use psb_c_cuda_vect_mod
#else
use psb_c_cuda_elg_mat_mod
#endif
implicit none
class(psb_c_cuda_elg_sparse_mat), intent(inout) :: a

7
cuda/impl/psb_c_cuda_elg_from_gpu.F90
Unescape Escape View File

@ -29,17 +29,12 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_elg_from_gpu(a,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use elldev_mod
use psb_vectordev_mod
use psb_c_cuda_elg_mat_mod, psb_protect_name => psb_c_cuda_elg_from_gpu
#else
use psb_c_cuda_elg_mat_mod
#endif
implicit none
class(psb_c_cuda_elg_sparse_mat), intent(inout) :: a
integer(psb_ipk_), intent(out) :: info
@ -48,7 +43,6 @@ subroutine psb_c_cuda_elg_from_gpu(a,info)
info = 0
#ifdef HAVE_SPGPU
if (.not.(c_associated(a%deviceMat))) then
call a%free()
return
@ -69,6 +63,5 @@ subroutine psb_c_cuda_elg_from_gpu(a,info)
if (info == 0) info = &
& readEllDevice(a%deviceMat,a%val,a%ja,pitch,a%irn,a%idiag)
call a%set_sync()
#endif
end subroutine psb_c_cuda_elg_from_gpu

5
cuda/impl/psb_c_cuda_elg_inner_vect_sv.F90
Unescape Escape View File

@ -29,17 +29,12 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_elg_inner_vect_sv(alpha,a,x,beta,y,info,trans)
use psb_base_mod
#ifdef HAVE_SPGPU
use elldev_mod
use psb_vectordev_mod
use psb_c_cuda_elg_mat_mod, psb_protect_name => psb_c_cuda_elg_inner_vect_sv
#else
use psb_c_cuda_elg_mat_mod
#endif
use psb_c_cuda_vect_mod
implicit none
class(psb_c_cuda_elg_sparse_mat), intent(in) :: a

2
cuda/impl/psb_c_cuda_elg_mold.F90
Unescape Escape View File

@ -28,8 +28,6 @@
! ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_elg_mold(a,b,info)
use psb_base_mod

7
cuda/impl/psb_c_cuda_elg_reallocate_nz.F90
Unescape Escape View File

@ -29,17 +29,12 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_elg_reallocate_nz(nz,a)
use psb_base_mod
#ifdef HAVE_SPGPU
use elldev_mod
use psb_vectordev_mod
use psb_c_cuda_elg_mat_mod, psb_protect_name => psb_c_cuda_elg_reallocate_nz
#else
use psb_c_cuda_elg_mat_mod
#endif
implicit none
integer(psb_ipk_), intent(in) :: nz
class(psb_c_cuda_elg_sparse_mat), intent(inout) :: a
@ -64,10 +59,8 @@ subroutine psb_c_cuda_elg_reallocate_nz(nz,a)
goto 9999
end if
#ifdef HAVE_SPGPU
call a%to_gpu(info,nzrm=nzrm)
if (info /= 0) goto 9999
#endif
call psb_erractionrestore(err_act)
return

7
cuda/impl/psb_c_cuda_elg_scal.F90
Unescape Escape View File

@ -29,17 +29,12 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_elg_scal(d,a,info,side)
use psb_base_mod
#ifdef HAVE_SPGPU
use elldev_mod
use psb_vectordev_mod
use psb_c_cuda_elg_mat_mod, psb_protect_name => psb_c_cuda_elg_scal
#else
use psb_c_cuda_elg_mat_mod
#endif
implicit none
class(psb_c_cuda_elg_sparse_mat), intent(inout) :: a
complex(psb_spk_), intent(in) :: d(:)
@ -63,10 +58,8 @@ subroutine psb_c_cuda_elg_scal(d,a,info,side)
call a%psb_c_ell_sparse_mat%scal(d,info,side)
if (info /= psb_success_) goto 9999
#ifdef HAVE_SPGPU
call a%to_gpu(info)
if (info /= 0) goto 9999
#endif
call psb_erractionrestore(err_act)
return

7
cuda/impl/psb_c_cuda_elg_scals.F90
Unescape Escape View File

@ -29,17 +29,12 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_elg_scals(d,a,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use elldev_mod
use psb_vectordev_mod
use psb_c_cuda_elg_mat_mod, psb_protect_name => psb_c_cuda_elg_scals
#else
use psb_c_cuda_elg_mat_mod
#endif
implicit none
class(psb_c_cuda_elg_sparse_mat), intent(inout) :: a
complex(psb_spk_), intent(in) :: d
@ -59,10 +54,8 @@ subroutine psb_c_cuda_elg_scals(d,a,info)
a%val(:,:) = a%val(:,:) * d
#ifdef HAVE_SPGPU
call a%to_gpu(info)
if (info /= 0) goto 9999
#endif
return

9
cuda/impl/psb_c_cuda_elg_to_gpu.F90
Unescape Escape View File

@ -29,30 +29,22 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_elg_to_gpu(a,info,nzrm)
use psb_base_mod
#ifdef HAVE_SPGPU
use elldev_mod
use psb_vectordev_mod
use psb_c_cuda_elg_mat_mod, psb_protect_name => psb_c_cuda_elg_to_gpu
#else
use psb_c_cuda_elg_mat_mod
#endif
implicit none
class(psb_c_cuda_elg_sparse_mat), intent(inout) :: a
integer(psb_ipk_), intent(out) :: info
integer(psb_ipk_), intent(in), optional :: nzrm
integer(psb_ipk_) :: m, nzm, n, pitch,maxrowsize, nzt
#ifdef HAVE_SPGPU
type(elldev_parms) :: gpu_parms
#endif
info = 0
#ifdef HAVE_SPGPU
if ((.not.allocated(a%val)).or.(.not.allocated(a%ja))) return
m = a%get_nrows()
@ -88,6 +80,5 @@ subroutine psb_c_cuda_elg_to_gpu(a,info,nzrm)
if (info == 0) info = &
& writeEllDevice(a%deviceMat,a%val,a%ja,size(a%ja,1),a%irn,a%idiag)
call a%set_sync()
#endif
end subroutine psb_c_cuda_elg_to_gpu

1
cuda/impl/psb_c_cuda_elg_trim.f90
Unescape Escape View File

@ -29,7 +29,6 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_elg_trim(a)
use psb_base_mod

10
cuda/impl/psb_c_cuda_elg_vect_mv.F90
Unescape Escape View File

@ -29,17 +29,12 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_elg_vect_mv(alpha,a,x,beta,y,info,trans)
use psb_base_mod
#ifdef HAVE_SPGPU
use elldev_mod
use psb_vectordev_mod
use psb_c_cuda_elg_mat_mod, psb_protect_name => psb_c_cuda_elg_vect_mv
#else
use psb_c_cuda_elg_mat_mod
#endif
use psb_c_cuda_vect_mod
implicit none
class(psb_c_cuda_elg_sparse_mat), intent(in) :: a
@ -71,7 +66,6 @@ subroutine psb_c_cuda_elg_vect_mv(alpha,a,x,beta,y,info,trans)
tra = (psb_toupper(trans_) == 'T').or.(psb_toupper(trans_)=='C')
#ifdef HAVE_SPGPU
if (tra) then
if (a%is_dev()) call a%sync()
if (.not.x%is_host()) call x%sync()
@ -116,10 +110,6 @@ subroutine psb_c_cuda_elg_vect_mv(alpha,a,x,beta,y,info,trans)
end select
end if
#else
if (a%is_dev()) call a%sync()
call a%psb_c_ell_sparse_mat%spmm(alpha,x,beta,y,info,trans)
#endif
if (info /= 0) goto 9999
call psb_erractionrestore(err_act)
return

10
cuda/impl/psb_c_cuda_hdiag_csmv.F90
Unescape Escape View File

@ -29,17 +29,12 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_hdiag_csmv(alpha,a,x,beta,y,info,trans)
use psb_base_mod
#ifdef HAVE_SPGPU
use hdiagdev_mod
use psb_vectordev_mod
use psb_c_cuda_hdiag_mat_mod, psb_protect_name => psb_c_cuda_hdiag_csmv
#else
use psb_c_cuda_hdiag_mat_mod
#endif
implicit none
class(psb_c_cuda_hdiag_sparse_mat), intent(in) :: a
complex(psb_spk_), intent(in) :: alpha, beta, x(:)
@ -94,7 +89,6 @@ subroutine psb_c_cuda_hdiag_csmv(alpha,a,x,beta,y,info,trans)
goto 9999
end if
#ifdef HAVE_SPGPU
if (tra) then
call a%psb_c_hdia_sparse_mat%spmm(alpha,x,beta,y,info,trans)
else
@ -121,9 +115,6 @@ subroutine psb_c_cuda_hdiag_csmv(alpha,a,x,beta,y,info,trans)
call freeMultiVecDevice(gpX)
call freeMultiVecDevice(gpY)
endif
#else
call a%psb_c_hdia_sparse_mat%spmm(alpha,x,beta,y,info,trans)
#endif
call psb_erractionrestore(err_act)
return
@ -132,5 +123,4 @@ subroutine psb_c_cuda_hdiag_csmv(alpha,a,x,beta,y,info,trans)
return
end subroutine psb_c_cuda_hdiag_csmv

1
cuda/impl/psb_c_cuda_hdiag_mold.F90
Unescape Escape View File

@ -29,7 +29,6 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_hdiag_mold(a,b,info)
use psb_base_mod

10
cuda/impl/psb_c_cuda_hdiag_to_gpu.F90
Unescape Escape View File

@ -29,29 +29,21 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_hdiag_to_gpu(a,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use hdiagdev_mod
use psb_vectordev_mod
use psb_c_cuda_hdiag_mat_mod, psb_protect_name => psb_c_cuda_hdiag_to_gpu
#else
use psb_c_cuda_hdiag_mat_mod
#endif
use iso_c_binding
implicit none
class(psb_c_cuda_hdiag_sparse_mat), intent(inout) :: a
integer(psb_ipk_), intent(out) :: info
integer(psb_ipk_) :: nr, nc, hacksize, hackcount, allocheight
#ifdef HAVE_SPGPU
type(hdiagdev_parms) :: gpu_parms
#endif
info = 0
#ifdef HAVE_SPGPU
nr = a%get_nrows()
nc = a%get_ncols()
hacksize = a%hackSize
@ -81,6 +73,4 @@ subroutine psb_c_cuda_hdiag_to_gpu(a,info)
if (info == 0) info = &
& writeHdiagDevice(a%deviceMat,a%val,a%diaOffsets,a%hackOffsets)
#endif
end subroutine psb_c_cuda_hdiag_to_gpu

9
cuda/impl/psb_c_cuda_hdiag_vect_mv.F90
Unescape Escape View File

@ -29,17 +29,12 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_hdiag_vect_mv(alpha,a,x,beta,y,info,trans)
use psb_base_mod
#ifdef HAVE_SPGPU
use hdiagdev_mod
use psb_vectordev_mod
use psb_c_cuda_hdiag_mat_mod, psb_protect_name => psb_c_cuda_hdiag_vect_mv
#else
use psb_c_cuda_hdiag_mat_mod
#endif
use psb_c_cuda_vect_mod
implicit none
class(psb_c_cuda_hdiag_sparse_mat), intent(in) :: a
@ -71,7 +66,6 @@ subroutine psb_c_cuda_hdiag_vect_mv(alpha,a,x,beta,y,info,trans)
tra = (psb_toupper(trans_) == 'T').or.(psb_toupper(trans_)=='C')
#ifdef HAVE_SPGPU
if (tra) then
if (.not.x%is_host()) call x%sync()
if (beta /= dzero) then
@ -112,9 +106,6 @@ subroutine psb_c_cuda_hdiag_vect_mv(alpha,a,x,beta,y,info,trans)
end select
end if
#else
call a%psb_c_hdia_sparse_mat%spmm(alpha,x,beta,y,info,trans)
#endif
if (info /= 0) goto 9999
call psb_erractionrestore(err_act)
return

9
cuda/impl/psb_c_cuda_hlg_allocate_mnnz.F90
Unescape Escape View File

@ -29,17 +29,12 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_hlg_allocate_mnnz(m,n,a,nz)
use psb_base_mod
#ifdef HAVE_SPGPU
use hlldev_mod
use psb_vectordev_mod
use psb_c_cuda_hlg_mat_mod, psb_protect_name => psb_c_cuda_hlg_allocate_mnnz
#else
use psb_c_cuda_hlg_mat_mod
#endif
implicit none
integer(psb_ipk_), intent(in) :: m,n
class(psb_c_cuda_hlg_sparse_mat), intent(inout) :: a
@ -47,19 +42,15 @@ subroutine psb_c_cuda_hlg_allocate_mnnz(m,n,a,nz)
Integer(psb_ipk_) :: err_act, info, nz_,ld
character(len=20) :: name='allocate_mnz'
logical, parameter :: debug=.false.
#ifdef HAVE_SPGPU
type(hlldev_parms) :: gpu_parms
#endif
call psb_erractionsave(err_act)
info = psb_success_
call a%psb_c_hll_sparse_mat%allocate(m,n,nz)
#ifdef HAVE_SPGPU
call a%to_gpu(info,nzrm=nz_)
if (info /= 0) goto 9999
#endif
call psb_erractionrestore(err_act)
return

10
cuda/impl/psb_c_cuda_hlg_csmm.F90
Unescape Escape View File

@ -29,17 +29,12 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_hlg_csmm(alpha,a,x,beta,y,info,trans)
use psb_base_mod
#ifdef HAVE_SPGPU
use hlldev_mod
use psb_vectordev_mod
use psb_c_cuda_hlg_mat_mod, psb_protect_name => psb_c_cuda_hlg_csmm
#else
use psb_c_cuda_hlg_mat_mod
#endif
implicit none
class(psb_c_cuda_hlg_sparse_mat), intent(in) :: a
complex(psb_spk_), intent(in) :: alpha, beta, x(:,:)
@ -93,7 +88,6 @@ subroutine psb_c_cuda_hlg_csmm(alpha,a,x,beta,y,info,trans)
end if
#ifdef HAVE_SPGPU
if (tra) then
call a%psb_c_hll_sparse_mat%spmm(alpha,x,beta,y,info,trans)
else
@ -118,9 +112,6 @@ subroutine psb_c_cuda_hlg_csmm(alpha,a,x,beta,y,info,trans)
call freeMultiVecDevice(gpX)
call freeMultiVecDevice(gpY)
endif
#else
call a%psb_c_hll_sparse_mat%spmm(alpha,x,beta,y,info,trans)
#endif
call psb_erractionrestore(err_act)
return
@ -128,5 +119,4 @@ subroutine psb_c_cuda_hlg_csmm(alpha,a,x,beta,y,info,trans)
return
end subroutine psb_c_cuda_hlg_csmm

9
cuda/impl/psb_c_cuda_hlg_csmv.F90
Unescape Escape View File

@ -29,17 +29,12 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_hlg_csmv(alpha,a,x,beta,y,info,trans)
use psb_base_mod
#ifdef HAVE_SPGPU
use hlldev_mod
use psb_vectordev_mod
use psb_c_cuda_hlg_mat_mod, psb_protect_name => psb_c_cuda_hlg_csmv
#else
use psb_c_cuda_hlg_mat_mod
#endif
implicit none
class(psb_c_cuda_hlg_sparse_mat), intent(in) :: a
complex(psb_spk_), intent(in) :: alpha, beta, x(:)
@ -94,7 +89,6 @@ subroutine psb_c_cuda_hlg_csmv(alpha,a,x,beta,y,info,trans)
goto 9999
end if
#ifdef HAVE_SPGPU
if (tra) then
call a%psb_c_hll_sparse_mat%spmm(alpha,x,beta,y,info,trans)
else
@ -121,9 +115,6 @@ subroutine psb_c_cuda_hlg_csmv(alpha,a,x,beta,y,info,trans)
call freeMultiVecDevice(gpX)
call freeMultiVecDevice(gpY)
endif
#else
call a%psb_c_hll_sparse_mat%spmm(alpha,x,beta,y,info,trans)
#endif
call psb_erractionrestore(err_act)
return

7
cuda/impl/psb_c_cuda_hlg_from_gpu.F90
Unescape Escape View File

@ -29,17 +29,12 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_hlg_from_gpu(a,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use hlldev_mod
use psb_vectordev_mod
use psb_c_cuda_hlg_mat_mod, psb_protect_name => psb_c_cuda_hlg_from_gpu
#else
use psb_c_cuda_hlg_mat_mod
#endif
implicit none
class(psb_c_cuda_hlg_sparse_mat), intent(inout) :: a
integer(psb_ipk_), intent(out) :: info
@ -48,7 +43,6 @@ subroutine psb_c_cuda_hlg_from_gpu(a,info)
info = 0
#ifdef HAVE_SPGPU
if (a%is_sync()) return
if (a%is_host()) return
if (.not.(c_associated(a%deviceMat))) then
@ -71,6 +65,5 @@ subroutine psb_c_cuda_hlg_from_gpu(a,info)
if (info == 0) info = &
& readHllDevice(a%deviceMat,a%val,a%ja,a%hkoffs,a%irn,a%idiag)
call a%set_sync()
#endif
end subroutine psb_c_cuda_hlg_from_gpu

7
cuda/impl/psb_c_cuda_hlg_inner_vect_sv.F90
Unescape Escape View File

@ -29,17 +29,12 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_hlg_inner_vect_sv(alpha,a,x,beta,y,info,trans)
use psb_base_mod
#ifdef HAVE_SPGPU
use hlldev_mod
use psb_vectordev_mod
use psb_c_cuda_hlg_mat_mod, psb_protect_name => psb_c_cuda_hlg_inner_vect_sv
#else
use psb_c_cuda_hlg_mat_mod
#endif
use psb_c_cuda_vect_mod
implicit none
class(psb_c_cuda_hlg_sparse_mat), intent(in) :: a
@ -69,11 +64,9 @@ subroutine psb_c_cuda_hlg_inner_vect_sv(alpha,a,x,beta,y,info,trans)
goto 9999
end if
call psb_erractionrestore(err_act)
return
9999 call psb_error_handler(err_act)
return

1
cuda/impl/psb_c_cuda_hlg_mold.F90
Unescape Escape View File

@ -29,7 +29,6 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_hlg_mold(a,b,info)
use psb_base_mod

7
cuda/impl/psb_c_cuda_hlg_reallocate_nz.F90
Unescape Escape View File

@ -29,17 +29,12 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_hlg_reallocate_nz(nz,a)
use psb_base_mod
#ifdef HAVE_SPGPU
use hlldev_mod
use psb_vectordev_mod
use psb_c_cuda_hlg_mat_mod, psb_protect_name => psb_c_cuda_hlg_reallocate_nz
#else
use psb_c_cuda_hlg_mat_mod
#endif
use iso_c_binding
implicit none
integer(psb_ipk_), intent(in) :: nz
@ -52,10 +47,8 @@ subroutine psb_c_cuda_hlg_reallocate_nz(nz,a)
call a%psb_c_hll_sparse_mat%reallocate(nz)
#ifdef HAVE_SPGPU
call a%to_gpu(info)
if (info /= 0) goto 9999
#endif
call psb_erractionrestore(err_act)
return

7
cuda/impl/psb_c_cuda_hlg_scal.F90
Unescape Escape View File

@ -29,17 +29,12 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_hlg_scal(d,a,info,side)
use psb_base_mod
#ifdef HAVE_SPGPU
use hlldev_mod
use psb_vectordev_mod
use psb_c_cuda_hlg_mat_mod, psb_protect_name => psb_c_cuda_hlg_scal
#else
use psb_c_cuda_hlg_mat_mod
#endif
implicit none
class(psb_c_cuda_hlg_sparse_mat), intent(inout) :: a
complex(psb_spk_), intent(in) :: d(:)
@ -60,10 +55,8 @@ subroutine psb_c_cuda_hlg_scal(d,a,info,side)
call a%psb_c_hll_sparse_mat%scal(d,info,side)
if (info /= psb_success_) goto 9999
#ifdef HAVE_SPGPU
call a%to_gpu(info)
if (info /= 0) goto 9999
#endif
call psb_erractionrestore(err_act)
return

7
cuda/impl/psb_c_cuda_hlg_scals.F90
Unescape Escape View File

@ -29,17 +29,12 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_hlg_scals(d,a,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use hlldev_mod
use psb_vectordev_mod
use psb_c_cuda_hlg_mat_mod, psb_protect_name => psb_c_cuda_hlg_scals
#else
use psb_c_cuda_hlg_mat_mod
#endif
use iso_c_binding
implicit none
class(psb_c_cuda_hlg_sparse_mat), intent(inout) :: a
@ -59,10 +54,8 @@ subroutine psb_c_cuda_hlg_scals(d,a,info)
call a%psb_c_hll_sparse_mat%scal(d,info)
if (info /= psb_success_) goto 9999
#ifdef HAVE_SPGPU
call a%to_gpu(info)
if (info /= 0) goto 9999
#endif
call psb_erractionrestore(err_act)
return

7
cuda/impl/psb_c_cuda_hlg_to_gpu.F90
Unescape Escape View File

@ -29,17 +29,12 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_hlg_to_gpu(a,info,nzrm)
use psb_base_mod
#ifdef HAVE_SPGPU
use hlldev_mod
use psb_vectordev_mod
use psb_c_cuda_hlg_mat_mod, psb_protect_name => psb_c_cuda_hlg_to_gpu
#else
use psb_c_cuda_hlg_mat_mod
#endif
use iso_c_binding
implicit none
class(psb_c_cuda_hlg_sparse_mat), intent(inout) :: a
@ -50,7 +45,6 @@ subroutine psb_c_cuda_hlg_to_gpu(a,info,nzrm)
info = 0
#ifdef HAVE_SPGPU
if ((.not.allocated(a%val)).or.(.not.allocated(a%ja))) return
n = a%get_nrows()
@ -63,6 +57,5 @@ subroutine psb_c_cuda_hlg_to_gpu(a,info,nzrm)
if (info == 0) info = &
& writehllDevice(a%deviceMat,a%val,a%ja,a%hkoffs,a%irn,a%idiag)
! if (info /= 0) goto 9999
#endif
end subroutine psb_c_cuda_hlg_to_gpu

10
cuda/impl/psb_c_cuda_hlg_vect_mv.F90
Unescape Escape View File

@ -29,17 +29,12 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_hlg_vect_mv(alpha,a,x,beta,y,info,trans)
use psb_base_mod
#ifdef HAVE_SPGPU
use hlldev_mod
use psb_vectordev_mod
use psb_c_cuda_hlg_mat_mod, psb_protect_name => psb_c_cuda_hlg_vect_mv
#else
use psb_c_cuda_hlg_mat_mod
#endif
use psb_c_cuda_vect_mod
implicit none
class(psb_c_cuda_hlg_sparse_mat), intent(in) :: a
@ -69,9 +64,7 @@ subroutine psb_c_cuda_hlg_vect_mv(alpha,a,x,beta,y,info,trans)
goto 9999
endif
tra = (psb_toupper(trans_) == 'T').or.(psb_toupper(trans_)=='C')
#ifdef HAVE_SPGPU
if (tra) then
if (.not.x%is_host()) call x%sync()
if (beta /= czero) then
@ -115,9 +108,6 @@ subroutine psb_c_cuda_hlg_vect_mv(alpha,a,x,beta,y,info,trans)
end select
end if
#else
call a%psb_c_hll_sparse_mat%spmm(alpha,x,beta,y,info,trans)
#endif
if (info /= 0) goto 9999
call psb_erractionrestore(err_act)
return

6
cuda/impl/psb_c_cuda_hybg_allocate_mnnz.F90
Unescape Escape View File

@ -33,12 +33,8 @@
subroutine psb_c_cuda_hybg_allocate_mnnz(m,n,a,nz)
use psb_base_mod
#ifdef HAVE_SPGPU
use cusparse_mod
use psb_c_cuda_hybg_mat_mod, psb_protect_name => psb_c_cuda_hybg_allocate_mnnz
#else
use psb_c_cuda_hybg_mat_mod
#endif
implicit none
integer(psb_ipk_), intent(in) :: m,n
class(psb_c_cuda_hybg_sparse_mat), intent(inout) :: a
@ -52,11 +48,9 @@ subroutine psb_c_cuda_hybg_allocate_mnnz(m,n,a,nz)
call a%psb_c_csr_sparse_mat%allocate(m,n,nz)
#ifdef HAVE_SPGPU
info = initFcusparse()
call a%to_gpu(info,nzrm=nz)
if (info /= 0) goto 9999
#endif
call psb_erractionrestore(err_act)
return

9
cuda/impl/psb_c_cuda_hybg_csmm.F90
Unescape Escape View File

@ -33,14 +33,10 @@
subroutine psb_c_cuda_hybg_csmm(alpha,a,x,beta,y,info,trans)
use psb_base_mod
#ifdef HAVE_SPGPU
use cusparse_mod
use elldev_mod
use psb_vectordev_mod
use psb_c_cuda_hybg_mat_mod, psb_protect_name => psb_c_cuda_hybg_csmm
#else
use psb_c_cuda_hybg_mat_mod
#endif
implicit none
class(psb_c_cuda_hybg_sparse_mat), intent(in) :: a
complex(psb_spk_), intent(in) :: alpha, beta, x(:,:)
@ -92,8 +88,6 @@ subroutine psb_c_cuda_hybg_csmm(alpha,a,x,beta,y,info,trans)
goto 9999
end if
#ifdef HAVE_SPGPU
if (tra) then
call a%psb_c_csr_sparse_mat%spmm(alpha,x,beta,y,info,trans)
else
@ -121,9 +115,6 @@ subroutine psb_c_cuda_hybg_csmm(alpha,a,x,beta,y,info,trans)
call freeMultiVecDevice(gpX)
call freeMultiVecDevice(gpY)
endif
#else
call a%psb_c_csr_sparse_mat%spmm(alpha,x,beta,y,info,trans)
#endif
call psb_erractionrestore(err_act)
return

8
cuda/impl/psb_c_cuda_hybg_csmv.F90
Unescape Escape View File

@ -33,14 +33,10 @@
subroutine psb_c_cuda_hybg_csmv(alpha,a,x,beta,y,info,trans)
use psb_base_mod
#ifdef HAVE_SPGPU
use cusparse_mod
use elldev_mod
use psb_vectordev_mod
use psb_c_cuda_hybg_mat_mod, psb_protect_name => psb_c_cuda_hybg_csmv
#else
use psb_c_cuda_hybg_mat_mod
#endif
implicit none
class(psb_c_cuda_hybg_sparse_mat), intent(in) :: a
complex(psb_spk_), intent(in) :: alpha, beta, x(:)
@ -95,7 +91,6 @@ subroutine psb_c_cuda_hybg_csmv(alpha,a,x,beta,y,info,trans)
goto 9999
end if
#ifdef HAVE_SPGPU
if (tra) then
call a%psb_c_csr_sparse_mat%spmm(alpha,x,beta,y,info,trans)
else
@ -123,9 +118,6 @@ subroutine psb_c_cuda_hybg_csmv(alpha,a,x,beta,y,info,trans)
call freeMultiVecDevice(gpX)
call freeMultiVecDevice(gpY)
endif
#else
call a%psb_c_csr_sparse_mat%spmm(alpha,x,beta,y,info,trans)
#endif
call psb_erractionrestore(err_act)
return

11
cuda/impl/psb_c_cuda_hybg_inner_vect_sv.F90
Unescape Escape View File

@ -33,13 +33,9 @@
subroutine psb_c_cuda_hybg_inner_vect_sv(alpha,a,x,beta,y,info,trans)
use psb_base_mod
#ifdef HAVE_SPGPU
use elldev_mod
use psb_vectordev_mod
use psb_c_cuda_hybg_mat_mod, psb_protect_name => psb_c_cuda_hybg_inner_vect_sv
#else
use psb_c_cuda_hybg_mat_mod
#endif
use psb_c_cuda_vect_mod
implicit none
class(psb_c_cuda_hybg_sparse_mat), intent(in) :: a
@ -76,7 +72,6 @@ subroutine psb_c_cuda_hybg_inner_vect_sv(alpha,a,x,beta,y,info,trans)
tra = (psb_toupper(trans_) == 'T').or.(psb_toupper(trans_)=='C')
#ifdef HAVE_SPGPU
if (tra.or.(beta/=czero)) then
call x%sync()
call y%sync()
@ -113,12 +108,6 @@ subroutine psb_c_cuda_hybg_inner_vect_sv(alpha,a,x,beta,y,info,trans)
call y%bld(ry)
end select
end if
#else
call x%sync()
call y%sync()
call a%psb_c_csr_sparse_mat%inner_spsm(alpha,x,beta,y,info,trans)
call y%set_host()
#endif
if (info /= psb_success_) then
info = psb_err_from_subroutine_
call psb_errpush(info,name, a_err='hybg_vect_sv')

6
cuda/impl/psb_c_cuda_hybg_reallocate_nz.F90
Unescape Escape View File

@ -33,12 +33,8 @@
subroutine psb_c_cuda_hybg_reallocate_nz(nz,a)
use psb_base_mod
#ifdef HAVE_SPGPU
use cusparse_mod
use psb_c_cuda_hybg_mat_mod, psb_protect_name => psb_c_cuda_hybg_reallocate_nz
#else
use psb_c_cuda_hybg_mat_mod
#endif
implicit none
integer(psb_ipk_), intent(in) :: nz
class(psb_c_cuda_hybg_sparse_mat), intent(inout) :: a
@ -55,10 +51,8 @@ subroutine psb_c_cuda_hybg_reallocate_nz(nz,a)
!
call a%psb_c_csr_sparse_mat%reallocate(nz)
#ifdef HAVE_SPGPU
call a%to_gpu(info,nzrm=nz)
if (info /= 0) goto 9999
#endif
call psb_erractionrestore(err_act)
return

6
cuda/impl/psb_c_cuda_hybg_scal.F90
Unescape Escape View File

@ -33,12 +33,8 @@
subroutine psb_c_cuda_hybg_scal(d,a,info,side)
use psb_base_mod
#ifdef HAVE_SPGPU
use cusparse_mod
use psb_c_cuda_hybg_mat_mod, psb_protect_name => psb_c_cuda_hybg_scal
#else
use psb_c_cuda_hybg_mat_mod
#endif
implicit none
class(psb_c_cuda_hybg_sparse_mat), intent(inout) :: a
complex(psb_spk_), intent(in) :: d(:)
@ -60,10 +56,8 @@ subroutine psb_c_cuda_hybg_scal(d,a,info,side)
call a%psb_c_csr_sparse_mat%scal(d,info,side=side)
if (info /= 0) goto 9999
#ifdef HAVE_SPGPU
call a%to_gpu(info)
if (info /= 0) goto 9999
#endif
call psb_erractionrestore(err_act)
return

6
cuda/impl/psb_c_cuda_hybg_scals.F90
Unescape Escape View File

@ -33,12 +33,8 @@
subroutine psb_c_cuda_hybg_scals(d,a,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use cusparse_mod
use psb_c_cuda_hybg_mat_mod, psb_protect_name => psb_c_cuda_hybg_scals
#else
use psb_c_cuda_hybg_mat_mod
#endif
implicit none
class(psb_c_cuda_hybg_sparse_mat), intent(inout) :: a
complex(psb_spk_), intent(in) :: d
@ -60,10 +56,8 @@ subroutine psb_c_cuda_hybg_scals(d,a,info)
if (info /= 0) goto 9999
#ifdef HAVE_SPGPU
call a%to_gpu(info)
if (info /= 0) goto 9999
#endif
call psb_erractionrestore(err_act)
return

6
cuda/impl/psb_c_cuda_hybg_to_gpu.F90
Unescape Escape View File

@ -33,12 +33,8 @@
subroutine psb_c_cuda_hybg_to_gpu(a,info,nzrm)
use psb_base_mod
#ifdef HAVE_SPGPU
use cusparse_mod
use psb_c_cuda_hybg_mat_mod, psb_protect_name => psb_c_cuda_hybg_to_gpu
#else
use psb_c_cuda_hybg_mat_mod
#endif
implicit none
class(psb_c_cuda_hybg_sparse_mat), intent(inout) :: a
integer(psb_ipk_), intent(out) :: info
@ -51,7 +47,6 @@ subroutine psb_c_cuda_hybg_to_gpu(a,info,nzrm)
info = 0
#ifdef HAVE_SPGPU
if ((.not.allocated(a%val)).or.(.not.allocated(a%ja))) return
m = a%get_nrows()
@ -148,7 +143,6 @@ subroutine psb_c_cuda_hybg_to_gpu(a,info,nzrm)
if (info /= 0) then
write(0,*) 'Error in HYBG_TO_GPU ',info
end if
#endif
end subroutine psb_c_cuda_hybg_to_gpu
#endif

9
cuda/impl/psb_c_cuda_hybg_vect_mv.F90
Unescape Escape View File

@ -33,14 +33,10 @@
subroutine psb_c_cuda_hybg_vect_mv(alpha,a,x,beta,y,info,trans)
use psb_base_mod
#ifdef HAVE_SPGPU
use cusparse_mod
use elldev_mod
use psb_vectordev_mod
use psb_c_cuda_hybg_mat_mod, psb_protect_name => psb_c_cuda_hybg_vect_mv
#else
use psb_c_cuda_hybg_mat_mod
#endif
use psb_c_cuda_vect_mod
implicit none
class(psb_c_cuda_hybg_sparse_mat), intent(in) :: a
@ -71,8 +67,6 @@ subroutine psb_c_cuda_hybg_vect_mv(alpha,a,x,beta,y,info,trans)
tra = (psb_toupper(trans_) == 'T').or.(psb_toupper(trans_)=='C')
#ifdef HAVE_SPGPU
if (tra) then
if (.not.x%is_host()) call x%sync()
if (beta /= czero) then
@ -112,9 +106,6 @@ subroutine psb_c_cuda_hybg_vect_mv(alpha,a,x,beta,y,info,trans)
call y%bld(ry)
end select
end if
#else
call a%psb_c_csr_sparse_mat%spmm(alpha,x,beta,y,info,trans)
#endif
if (info /= 0) goto 9999
call psb_erractionrestore(err_act)
return

7
cuda/impl/psb_c_cuda_mv_csrg_from_coo.F90
Unescape Escape View File

@ -29,16 +29,11 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_mv_csrg_from_coo(a,b,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use cusparse_mod
use psb_c_cuda_csrg_mat_mod, psb_protect_name => psb_c_cuda_mv_csrg_from_coo
#else
use psb_c_cuda_csrg_mat_mod
#endif
implicit none
class(psb_c_cuda_csrg_sparse_mat), intent(inout) :: a
@ -51,9 +46,7 @@ subroutine psb_c_cuda_mv_csrg_from_coo(a,b,info)
call a%psb_c_csr_sparse_mat%mv_from_coo(b,info)
if (info /= 0) goto 9999
#ifdef HAVE_SPGPU
call a%to_gpu(info)
#endif
if (info /= 0) goto 9999
return

7
cuda/impl/psb_c_cuda_mv_csrg_from_fmt.F90
Unescape Escape View File

@ -29,16 +29,11 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_mv_csrg_from_fmt(a,b,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use cusparse_mod
use psb_c_cuda_csrg_mat_mod, psb_protect_name => psb_c_cuda_mv_csrg_from_fmt
#else
use psb_c_cuda_csrg_mat_mod
#endif
implicit none
class(psb_c_cuda_csrg_sparse_mat), intent(inout) :: a
@ -55,9 +50,7 @@ subroutine psb_c_cuda_mv_csrg_from_fmt(a,b,info)
class default
call a%psb_c_csr_sparse_mat%mv_from_fmt(b,info)
if (info /= 0) return
#ifdef HAVE_SPGPU
call a%to_gpu(info)
#endif
end select
end subroutine psb_c_cuda_mv_csrg_from_fmt

5
cuda/impl/psb_c_cuda_mv_diag_from_coo.F90
Unescape Escape View File

@ -29,17 +29,12 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_mv_diag_from_coo(a,b,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use diagdev_mod
use psb_vectordev_mod
use psb_c_cuda_diag_mat_mod, psb_protect_name => psb_c_cuda_mv_diag_from_coo
#else
use psb_c_cuda_diag_mat_mod
#endif
implicit none

6
cuda/impl/psb_c_cuda_mv_elg_from_coo.F90
Unescape Escape View File

@ -29,17 +29,12 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_mv_elg_from_coo(a,b,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use elldev_mod
use psb_vectordev_mod
use psb_c_cuda_elg_mat_mod, psb_protect_name => psb_c_cuda_mv_elg_from_coo
#else
use psb_c_cuda_elg_mat_mod
#endif
implicit none
class(psb_c_cuda_elg_sparse_mat), intent(inout) :: a
@ -57,5 +52,4 @@ subroutine psb_c_cuda_mv_elg_from_coo(a,b,info)
return
end subroutine psb_c_cuda_mv_elg_from_coo

13
cuda/impl/psb_c_cuda_mv_elg_from_fmt.F90
Unescape Escape View File

@ -29,17 +29,12 @@
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_mv_elg_from_fmt(a,b,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use elldev_mod
use psb_vectordev_mod
use psb_c_cuda_elg_mat_mod, psb_protect_name => psb_c_cuda_mv_elg_from_fmt
#else
use psb_c_cuda_elg_mat_mod
#endif
implicit none
class(psb_c_cuda_elg_sparse_mat), intent(inout) :: a
@ -49,9 +44,7 @@ subroutine psb_c_cuda_mv_elg_from_fmt(a,b,info)
!locals
type(psb_c_coo_sparse_mat) :: tmp
Integer(Psb_ipk_) :: nza, nr, i,j,irw, idl,err_act, nc, ld, nzm, m
#ifdef HAVE_SPGPU
type(elldev_parms) :: gpu_parms
#endif
info = psb_success_
@ -65,13 +58,9 @@ subroutine psb_c_cuda_mv_elg_from_fmt(a,b,info)
m = b%get_nrows()
nc = b%get_ncols()
nza = b%get_nzeros()
#ifdef HAVE_SPGPU
gpu_parms = FgetEllDeviceParams(m,nzm,nza,nc,spgpu_type_double,1)
ld = gpu_parms%pitch
nzm = gpu_parms%maxRowSize
#else
ld = m
#endif
a%psb_c_base_sparse_mat = b%psb_c_base_sparse_mat
call move_alloc(b%irn, a%irn)
call move_alloc(b%idiag, a%idiag)
@ -87,9 +76,7 @@ subroutine psb_c_cuda_mv_elg_from_fmt(a,b,info)
end if
a%nzt = nza
call b%free()
#ifdef HAVE_SPGPU
call a%to_gpu(info)
#endif
class default
call b%mv_to_coo(tmp,info)

10
cuda/impl/psb_c_cuda_mv_hdiag_from_coo.F90
Unescape Escape View File

@ -28,19 +28,13 @@
! ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_mv_hdiag_from_coo(a,b,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use hdiagdev_mod
use psb_vectordev_mod
use psb_c_cuda_hdiag_mat_mod, psb_protect_name => psb_c_cuda_mv_hdiag_from_coo
use psb_cuda_env_mod
#else
use psb_c_cuda_hdiag_mat_mod
#endif
implicit none
@ -54,16 +48,12 @@ subroutine psb_c_cuda_mv_hdiag_from_coo(a,b,info)
info = psb_success_
#ifdef HAVE_SPGPU
a%hacksize = psb_cuda_WarpSize()
#endif
call a%psb_c_hdia_sparse_mat%mv_from_coo(b,info)
#ifdef HAVE_SPGPU
call a%to_gpu(info)
if (info /= 0) goto 9999
#endif
return

6
cuda/impl/psb_c_cuda_mv_hlg_from_coo.F90
Unescape Escape View File

@ -28,19 +28,13 @@
! ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_mv_hlg_from_coo(a,b,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use hlldev_mod
use psb_vectordev_mod
use psb_cuda_env_mod
use psb_c_cuda_hlg_mat_mod, psb_protect_name => psb_c_cuda_mv_hlg_from_coo
#else
use psb_c_cuda_hlg_mat_mod
#endif
implicit none
class(psb_c_cuda_hlg_sparse_mat), intent(inout) :: a

6
cuda/impl/psb_c_cuda_mv_hlg_from_fmt.F90
Unescape Escape View File

@ -28,18 +28,12 @@
! ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
! POSSIBILITY OF SUCH DAMAGE.
!
subroutine psb_c_cuda_mv_hlg_from_fmt(a,b,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use hlldev_mod
use psb_vectordev_mod
use psb_c_cuda_hlg_mat_mod, psb_protect_name => psb_c_cuda_mv_hlg_from_fmt
#else
use psb_c_cuda_hlg_mat_mod
#endif
implicit none
class(psb_c_cuda_hlg_sparse_mat), intent(inout) :: a

6
cuda/impl/psb_c_cuda_mv_hybg_from_coo.F90
Unescape Escape View File

@ -33,12 +33,8 @@
subroutine psb_c_cuda_mv_hybg_from_coo(a,b,info)
use psb_base_mod
#ifdef HAVE_SPGPU
use cusparse_mod
use psb_c_cuda_hybg_mat_mod, psb_protect_name => psb_c_cuda_mv_hybg_from_coo
#else
use psb_c_cuda_hybg_mat_mod
#endif
implicit none
class(psb_c_cuda_hybg_sparse_mat), intent(inout) :: a
@ -50,10 +46,8 @@ subroutine psb_c_cuda_mv_hybg_from_coo(a,b,info)
call a%psb_c_csr_sparse_mat%mv_from_coo(b,info)
if (info /= 0) goto 9999
#ifdef HAVE_SPGPU
call a%to_gpu(info)
if (info /= 0) goto 9999
#endif
return

Some files were not shown because too many files have changed in this diff Show More

Write Preview
Loading…
Cancel
Save