ssymm.cu

Nividia提供的CUDA的BLAS库源码

               (params.upper << 1) | params.lside);

    cudaStat = cudaGetLastError(); /* clear error status */
    if (usePureHwStepper) {
        ssymm_hw[funcIdx]<<<ctaDimsHw,CUBLAS_SSYMM_THREAD_COUNT>>>(params);
    } else {
        ssymm_sw[funcIdx]<<<ctaDimsSw,CUBLAS_SSYMM_THREAD_COUNT>>>(params);
    }
    cudaStat = cudaGetLastError(); /* check for launch error */
    if (cudaStat != cudaSuccess) {
        cublasSetError (ctx, CUBLAS_STATUS_EXECUTION_FAILED);
    }
}

__global__ void ssymm_main_hw_lo_right (struct cublasSsymmParams parms)
{
#undef  USE_MIXED_STEPPER
#undef  UPPER
#undef  LSIDE
#undef  FAST_IMUL
#undef  FULL_TILES_ONLY
#define USE_MIXED_STEPPER 0
#define UPPER             0
#define LSIDE             0
#define FAST_IMUL         0
#define FULL_TILES_ONLY   0
#include "ssymm.h"
}

__global__ void ssymm_main_hw_lo_left (struct cublasSsymmParams parms)
{
#undef  USE_MIXED_STEPPER
#undef  UPPER
#undef  LSIDE
#undef  FAST_IMUL
#undef  FULL_TILES_ONLY
#define USE_MIXED_STEPPER 0
#define UPPER             0
#define LSIDE             1
#define FAST_IMUL         0
#define FULL_TILES_ONLY   0
#include "ssymm.h"
}

__global__ void ssymm_main_hw_up_right (struct cublasSsymmParams parms)
{
#undef  USE_MIXED_STEPPER
#undef  UPPER
#undef  LSIDE
#undef  FAST_IMUL
#undef  FULL_TILES_ONLY
#define USE_MIXED_STEPPER 0
#define UPPER             1
#define LSIDE             0
#define FAST_IMUL         0
#define FULL_TILES_ONLY   0
#include "ssymm.h"
}

__global__ void ssymm_main_hw_up_left (struct cublasSsymmParams parms)
{
#undef  USE_MIXED_STEPPER
#undef  UPPER
#undef  LSIDE
#undef  FAST_IMUL
#undef  FULL_TILES_ONLY
#define USE_MIXED_STEPPER 0
#define UPPER             1
#define LSIDE             1
#define FAST_IMUL         0
#define FULL_TILES_ONLY   0
#include "ssymm.h"
}

__global__ void ssymm_main_sw_lo_right (struct cublasSsymmParams parms)
{
#undef  USE_MIXED_STEPPER
#undef  UPPER
#undef  LSIDE
#undef  FAST_IMUL
#undef  FULL_TILES_ONLY
#define USE_MIXED_STEPPER 1
#define UPPER             0
#define LSIDE             0
#define FAST_IMUL         0
#define FULL_TILES_ONLY   0
#include "ssymm.h"
}

__global__ void ssymm_main_sw_lo_left (struct cublasSsymmParams parms)
{
#undef  USE_MIXED_STEPPER
#undef  UPPER
#undef  LSIDE
#undef  FAST_IMUL
#undef  FULL_TILES_ONLY
#define USE_MIXED_STEPPER 1
#define UPPER             0
#define LSIDE             1
#define FAST_IMUL         0
#define FULL_TILES_ONLY   0
#include "ssymm.h"
}

__global__ void ssymm_main_sw_up_right (struct cublasSsymmParams parms)
{
#undef  USE_MIXED_STEPPER
#undef  UPPER
#undef  LSIDE
#undef  FAST_IMUL
#undef  FULL_TILES_ONLY
#define USE_MIXED_STEPPER 1
#define UPPER             1
#define LSIDE             0
#define FAST_IMUL         0
#define FULL_TILES_ONLY   0
#include "ssymm.h"
}

__global__ void ssymm_main_sw_up_left (struct cublasSsymmParams parms)
{
#undef  USE_MIXED_STEPPER
#undef  UPPER
#undef  LSIDE
#undef  FAST_IMUL
#undef  FULL_TILES_ONLY
#define USE_MIXED_STEPPER 1
#define UPPER             1
#define LSIDE             1
#define FAST_IMUL         0
#define FULL_TILES_ONLY   0
#include "ssymm.h"
}

__global__ void ssymm_main_hw_lo_right_fulltile (struct cublasSsymmParams parms)
{
#undef  USE_MIXED_STEPPER
#undef  UPPER
#undef  LSIDE
#undef  FAST_IMUL
#undef  FULL_TILES_ONLY
#define USE_MIXED_STEPPER 0
#define UPPER             0
#define LSIDE             0
#define FAST_IMUL         0
#define FULL_TILES_ONLY   1
#include "ssymm.h"
}

__global__ void ssymm_main_hw_lo_left_fulltile (struct cublasSsymmParams parms)
{
#undef  USE_MIXED_STEPPER
#undef  UPPER
#undef  LSIDE
#undef  FAST_IMUL
#undef  FULL_TILES_ONLY
#define USE_MIXED_STEPPER 0
#define UPPER             0
#define LSIDE             1
#define FAST_IMUL         0
#define FULL_TILES_ONLY   1
#include "ssymm.h"
}

__global__ void ssymm_main_hw_up_right_fulltile (struct cublasSsymmParams parms)
{
#undef  USE_MIXED_STEPPER
#undef  UPPER
#undef  LSIDE
#undef  FAST_IMUL
#undef  FULL_TILES_ONLY
#define USE_MIXED_STEPPER 0
#define UPPER             1
#define LSIDE             0
#define FAST_IMUL         0
#define FULL_TILES_ONLY   1
#include "ssymm.h"
}

__global__ void ssymm_main_hw_up_left_fulltile (struct cublasSsymmParams parms)
{
#undef  USE_MIXED_STEPPER
#undef  UPPER
#undef  LSIDE
#undef  FAST_IMUL
#undef  FULL_TILES_ONLY
#define USE_MIXED_STEPPER 0
#define UPPER             1
#define LSIDE             1
#define FAST_IMUL         0
#define FULL_TILES_ONLY   1
#include "ssymm.h"
}

__global__ void ssymm_main_sw_lo_right_fulltile (struct cublasSsymmParams parms)
{
#undef  USE_MIXED_STEPPER
#undef  UPPER
#undef  LSIDE
#undef  FAST_IMUL
#undef  FULL_TILES_ONLY
#define USE_MIXED_STEPPER 1
#define UPPER             0
#define LSIDE             0
#define FAST_IMUL         0
#define FULL_TILES_ONLY   1
#include "ssymm.h"
}

__global__ void ssymm_main_sw_lo_left_fulltile (struct cublasSsymmParams parms)
{
#undef  USE_MIXED_STEPPER
#undef  UPPER
#undef  LSIDE
#undef  FAST_IMUL
#undef  FULL_TILES_ONLY
#define USE_MIXED_STEPPER 1
#define UPPER             0
#define LSIDE             1
#define FAST_IMUL         0
#define FULL_TILES_ONLY   1
#include "ssymm.h"
}

__global__ void ssymm_main_sw_up_right_fulltile (struct cublasSsymmParams parms)
{
#undef  USE_MIXED_STEPPER
#undef  UPPER
#undef  LSIDE
#undef  FAST_IMUL
#undef  FULL_TILES_ONLY
#define USE_MIXED_STEPPER 1
#define UPPER             1
#define LSIDE             0
#define FAST_IMUL         0
#define FULL_TILES_ONLY   1
#include "ssymm.h"
}

__global__ void ssymm_main_sw_up_left_fulltile (struct cublasSsymmParams parms)
{
#undef  USE_MIXED_STEPPER
#undef  UPPER
#undef  LSIDE
#undef  FAST_IMUL
#undef  FULL_TILES_ONLY
#define USE_MIXED_STEPPER 1
#define UPPER             1
#define LSIDE             1
#define FAST_IMUL         0
#define FULL_TILES_ONLY   1
#include "ssymm.h"
}