cublas_impl_kernels.h

Go to the documentation of this file.

/*
 *  SPDX-FileCopyrightText: Copyright 2021, Siavash Ameli <sameli@berkeley.edu>
 *  SPDX-License-Identifier: BSD-3-Clause
 *  SPDX-FileType: SOURCE
 *
 *  This program is free software: you can redistribute it and/or modify it
 *  under the terms of the license found in the LICENSE.txt file in the root
 *  directory of this source tree.
 */
 
#ifndef _CU_BASIC_ALGEBRA_CUBLAS_IMPL_KERNELS_H_
#define _CU_BASIC_ALGEBRA_CUBLAS_IMPL_KERNELS_H_
 
// =======
// Headers
// =======
 
// Restrict qualifier
#if defined(_MSC_VER)
    #define RESTRICT __restrict
#elif defined(__INTEL_COMPILER)
    #define RESTRICT __restrict
#elif defined(__CUDA__) || defined(__GNUC__) || defined(__clang__)
    #define RESTRICT __restrict__
#else
    #define RESTRICT
#endif
 
 
// ==================
// cublas impl kernel
// ==================
 
 
namespace cublas_impl_kernels
{
    // cublasTgemv kernel
    template <
        typename DataType, typename ComputeType, unsigned int block_size>
    __global__ void cublasTgemv_kernel(
            const bool trans,
            const int m,
            const int n,
            const DataType alpha,
            const DataType* RESTRICT A,
            const int lda,
            const DataType* RESTRICT x,
            const int incx,
            const DataType beta,
            DataType* RESTRICT y,
            const int incy);
 
    // cublasTcopy kernel
    template <typename DataType>
    __global__ void cublasTcopy_kernel(
            const int n,
            const DataType* RESTRICT x,
            const int incx,
            DataType* RESTRICT y,
            const int incy);
 
    // cublasTaxpy kernel
    template <typename DataType>
    __global__ void cublasTaxpy_kernel(
            const int n,
            const DataType alpha,
            const DataType* RESTRICT x,
            const int incx,
            DataType* RESTRICT y,
            const int incy);
 
    // cublasTdot kernel
    template <
        typename DataType, typename ComputeType, unsigned int block_size>
    __global__ void cublasTdot_kernel(
            const int n,
            const DataType* RESTRICT x,
            const int incx,
            const DataType* RESTRICT y,
            const int incy,
            ComputeType* RESTRICT result);
 
    // cublasTnrm2 kernel
    template <
        typename DataType, typename ComputeType, unsigned int block_size>
    __global__ void cublasTnrm2_kernel(
            const int n,
            const DataType* RESTRICT x,
            const int incx,
            ComputeType* RESTRICT result);
 
    // cublasTscal kernel
    template <typename DataType>
    __global__ void cublasTscal_kernel(
            const int n,
            const DataType alpha,
            DataType* RESTRICT x,
            const int incx);
}
 
#endif  // _CU_BASIC_ALGEBRA_CUBLAS_IMPL_KERNELS_H_