ArrayImplementations_GPU.cu

Go to the documentation of this file.

#include <assert.h>
#include <basic_gpu.h>
#include <arrayfunctions_gpu.h>

__global__
void ArrayCopy1D_kernel(  const double* x,  
                          double*       y,  
                          int n              )
{
    int tx = threadIdx.x + (blockIdx.x * blockDim.x);
    if (tx < n) y[tx] = x[tx];
    return;
}

__global__
void ArraySetValue_kernel(  double* x,      
                            int     n,      
                            double  value    )
{
    int tx = threadIdx.x + (blockIdx.x * blockDim.x);
    if (tx < n) x[tx] = value;
    return;
}

__global__
void ArrayAXPY_kernel(  const double* x,  
                        double        a,  
                        double*       y,  
                        int           n    )
{
    int tx = threadIdx.x + (blockIdx.x * blockDim.x);
    if (tx < n) y[tx] += a*x[tx];
    return;
}

__global__
void ArrayBlockMultiply_kernel( double*       x,  
                                const double* a,  
                                int           n,  
                                int           bs  )
{
    int tx = threadIdx.x + (blockIdx.x * blockDim.x);
    if (tx < n) {
        for (int i = 0; i < bs; i++) x[tx*bs + i] *= a[tx];
    }
}

__global__
void ArrayCopy1DNewScheme_kernel( const double* __restrict__ src,   
                                  double*       __restrict__ dest,  
                                  int           npoints,            
                                  int           nvars                )
{
    int p = blockDim.x * blockIdx.x + threadIdx.x;
    if (p < npoints) {
        for (int v=0; v<nvars; v++) {
            dest[p+v*npoints] = src[p*nvars+v];
        }
    }
    return;
}

void gpuSetDevice(int device )
{
    cudaSetDevice(device);
    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess) {
      fprintf(stderr,"Error in gpuSetDevice(): device=%d error message=\"%s\"\n", device, cudaGetErrorString(err));
    }
}

void gpuMemcpy( void*               dest, 
                const void*         src,  
                size_t              count,
                enum gpuMemcpyKind  kind   )
{
    switch (kind) {
        case gpuMemcpyHostToDevice:
            checkCuda( cudaMemcpy(dest, src, count, cudaMemcpyHostToDevice) );
            break;
        case gpuMemcpyDeviceToHost:
            checkCuda( cudaMemcpy(dest, src, count, cudaMemcpyDeviceToHost) );
            break;
        case gpuMemcpyDeviceToDevice:
            checkCuda( cudaMemcpy(dest, src, count, cudaMemcpyDeviceToDevice) );
            break;
        default:
            fprintf(stderr, "Error: invalid gpuMemcpyKind: %d\n", kind);
            assert(0);
            break;
    }
    return;
}

void gpuMalloc( void**  devPtr, 
                size_t  size     )
{
    cudaMalloc(devPtr, size);
    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess) {
      fprintf(  stderr,"Error in gpuMalloc(): size=%d, error message=\"%s\"\n", size, 
                cudaGetErrorString(err) );
    }
    return;
}

void gpuMemset( void*   devPtr, 
                int     value,  
                size_t  count    )
{
    checkCuda( cudaMemset(devPtr, value, count) );
    return;
}

void gpuFree(void *devPtr )
{
    checkCuda( cudaFree(devPtr) );
    return;
}

void gpuArrayCopy1D(  const double* x,  
                      double*       y,  
                      int           n    )
{
    int nblocks = (n - 1) / GPU_THREADS_PER_BLOCK + 1;
    ArrayCopy1D_kernel<<<nblocks, GPU_THREADS_PER_BLOCK>>>(x, y, n);
    cudaDeviceSynchronize();
    return;
}

void gpuArraySetValue(  double* devPtr, 
                        int     n,      
                        double  value    )
{
    int nblocks = (n - 1) / GPU_THREADS_PER_BLOCK + 1;
    ArraySetValue_kernel<<<nblocks, GPU_THREADS_PER_BLOCK>>>(devPtr, n, value);
    cudaDeviceSynchronize();
    return;
}

void gpuArrayAXPY(  const double*   x, 
                    double          a, 
                    double*         y, 
                    int             n   )
{
    int nblocks = (n - 1) / GPU_THREADS_PER_BLOCK + 1;
    ArrayAXPY_kernel<<<nblocks, GPU_THREADS_PER_BLOCK>>>(x, a, y, n);
    cudaDeviceSynchronize();
    return;
}

void gpuArrayBlockMultiply( double*       x,  
                            const double* a,  
                            int           n,  
                            int           bs   )
{
    int nblocks = (n - 1) / GPU_THREADS_PER_BLOCK + 1;
    ArrayBlockMultiply_kernel<<<nblocks, GPU_THREADS_PER_BLOCK>>>(x, a, n, bs);
    cudaDeviceSynchronize();
    return;
}

double gpuArraySumSquarenD(int nvars,  
                           int ndims,  
                           int *dim,   
                           int ghosts, 
                           int *index, 
                           double *x   
                          )
{
    double sum = 0;
    printf("gpuArraySumSquarenD hasn't been implemented, yet.\n");
    exit(0);
    return (sum);
}

void gpuArrayCopy1DNewScheme( const double* src,      
                              double*       dest,     
                              int           npoints,  
                              int           nvars      )
{
    int nblocks = (npoints-1) / GPU_THREADS_PER_BLOCK + 1;
    ArrayCopy1DNewScheme_kernel<<<nblocks, GPU_THREADS_PER_BLOCK>>>(src, dest, npoints, nvars);
    cudaDeviceSynchronize();
    return;
}

void gpuArrayCheckEqual( const char*   msg,      
                         const double* var_adj,  
                         const double* var_sep,  
                         int           npoints    )
{
    double *h_var_adj = (double *) malloc(npoints*sizeof(double));
    double *h_var_sep = (double *) malloc(npoints*sizeof(double));

    gpuMemcpy(h_var_adj, var_adj, npoints*sizeof(double), gpuMemcpyDeviceToHost);
    gpuMemcpy(h_var_sep, var_sep, npoints*sizeof(double), gpuMemcpyDeviceToHost);

    double max_err = 0.0;
    for (int j=0; j<npoints; j++) {
        if (h_var_sep[j] != h_var_adj[j]) {
            max_err = max(max_err, fabs(h_var_sep[j]-h_var_adj[j]));
        }
    }

    free(h_var_adj);
    free(h_var_sep);

    if (max_err > 1e-10) {
        printf("gpuArrayCheckEqual: %-30s max_err = %e\n", msg, max_err);
        exit(0);
    }
    return;
}