fast_sgemm.cu

来自「Nividia提供的CUDA的BLAS库源码」· CU 代码 · 共 741 行 · 第 1/2 页

#undef  TRANSA
#undef  TRANSB
#undef  FAST_IMUL
#undef  USE_TEX
#define FULL_TILES_ONLY   0
#define USE_MIXED_STEPPER 1
#define TRANSA            1
#define TRANSB            0
#define FAST_IMUL         1
#define USE_TEX           1
#include "sgemm.h"
}

__global__ void fast_sgemm_tex_main_sw_ta_tb (struct cublasSgemmParams parms)
{
#undef  FULL_TILES_ONLY
#undef  USE_MIXED_STEPPER
#undef  TRANSA
#undef  TRANSB
#undef  FAST_IMUL
#undef  USE_TEX
#define FULL_TILES_ONLY   0
#define USE_MIXED_STEPPER 1
#define TRANSA            1
#define TRANSB            1
#define FAST_IMUL         1
#define USE_TEX           1
#include "sgemm.h"
}

__global__ void fast_sgemm_tex_main_sw_na_nb_fulltile (struct cublasSgemmParams parms)
{
#undef  FULL_TILES_ONLY
#undef  USE_MIXED_STEPPER
#undef  TRANSA
#undef  TRANSB
#undef  FAST_IMUL
#undef  USE_TEX
#define FULL_TILES_ONLY   1
#define USE_MIXED_STEPPER 1
#define TRANSA            0
#define TRANSB            0
#define FAST_IMUL         1
#define USE_TEX           1
#include "sgemm.h"
}

__global__ void fast_sgemm_tex_main_sw_na_tb_fulltile (struct cublasSgemmParams parms)
{
#undef  FULL_TILES_ONLY
#undef  USE_MIXED_STEPPER
#undef  TRANSA
#undef  TRANSB
#undef  FAST_IMUL
#undef  USE_TEX
#define FULL_TILES_ONLY   1
#define USE_MIXED_STEPPER 1
#define TRANSA            0
#define TRANSB            1
#define FAST_IMUL         1
#define USE_TEX           1
#include "sgemm.h"
}

__global__ void fast_sgemm_tex_main_sw_ta_nb_fulltile (struct cublasSgemmParams parms)
{
#undef  FULL_TILES_ONLY
#undef  USE_MIXED_STEPPER
#undef  TRANSA
#undef  TRANSB
#undef  FAST_IMUL
#undef  USE_TEX
#define FULL_TILES_ONLY   1
#define USE_MIXED_STEPPER 1
#define TRANSA            1
#define TRANSB            0
#define FAST_IMUL         1
#define USE_TEX           1
#include "sgemm.h"
}

__global__ void fast_sgemm_tex_main_sw_ta_tb_fulltile (struct cublasSgemmParams parms)
{
#undef  FULL_TILES_ONLY
#undef  USE_MIXED_STEPPER
#undef  TRANSA
#undef  TRANSB
#undef  FAST_IMUL
#undef  USE_TEX
#define FULL_TILES_ONLY   1
#define USE_MIXED_STEPPER 1
#define TRANSA            1
#define TRANSB            1
#define FAST_IMUL         1
#define USE_TEX           1
#include "sgemm.h"
}

__global__ void fast_sgemm_gld_main_hw_na_nb (struct cublasSgemmParams parms)
{
#undef  FULL_TILES_ONLY
#undef  USE_MIXED_STEPPER
#undef  TRANSA
#undef  TRANSB
#undef  FAST_IMUL
#undef  USE_TEX
#define FULL_TILES_ONLY   0
#define USE_MIXED_STEPPER 0
#define TRANSA            0
#define TRANSB            0
#define FAST_IMUL         1
#define USE_TEX           0
#include "sgemm.h"
}

__global__ void fast_sgemm_gld_main_hw_na_tb (struct cublasSgemmParams parms)
{
#undef  FULL_TILES_ONLY
#undef  USE_MIXED_STEPPER
#undef  TRANSA
#undef  TRANSB
#undef  FAST_IMUL
#undef  USE_TEX
#define FULL_TILES_ONLY   0
#define USE_MIXED_STEPPER 0
#define TRANSA            0
#define TRANSB            1
#define FAST_IMUL         1
#define USE_TEX           0
#include "sgemm.h"
}

__global__ void fast_sgemm_gld_main_hw_ta_nb (struct cublasSgemmParams parms)
{
#undef  FULL_TILES_ONLY
#undef  USE_MIXED_STEPPER
#undef  TRANSA
#undef  TRANSB
#undef  FAST_IMUL
#undef  USE_TEX
#define FULL_TILES_ONLY   0
#define USE_MIXED_STEPPER 0
#define TRANSA            1
#define TRANSB            0
#define FAST_IMUL         1
#define USE_TEX           0
#include "sgemm.h"
}

__global__ void fast_sgemm_gld_main_hw_ta_tb (struct cublasSgemmParams parms)
{
#undef  FULL_TILES_ONLY
#undef  USE_MIXED_STEPPER
#undef  TRANSA
#undef  TRANSB
#undef  FAST_IMUL
#undef  USE_TEX
#define FULL_TILES_ONLY   0
#define USE_MIXED_STEPPER 0
#define TRANSA            1
#define TRANSB            1
#define FAST_IMUL         1
#define USE_TEX           0
#include "sgemm.h"
}

__global__ void fast_sgemm_gld_main_hw_na_nb_fulltile (struct cublasSgemmParams parms)
{
#undef  FULL_TILES_ONLY
#undef  USE_MIXED_STEPPER
#undef  TRANSA
#undef  TRANSB
#undef  FAST_IMUL
#undef  USE_TEX
#define FULL_TILES_ONLY   1
#define USE_MIXED_STEPPER 0
#define TRANSA            0
#define TRANSB            0
#define FAST_IMUL         1
#define USE_TEX           0
#include "sgemm.h"
}

__global__ void fast_sgemm_gld_main_hw_na_tb_fulltile (struct cublasSgemmParams parms)
{
#undef  FULL_TILES_ONLY
#undef  USE_MIXED_STEPPER
#undef  TRANSA
#undef  TRANSB
#undef  FAST_IMUL
#undef  USE_TEX
#define FULL_TILES_ONLY   1
#define USE_MIXED_STEPPER 0
#define TRANSA            0
#define TRANSB            1
#define FAST_IMUL         1
#define USE_TEX           0
#include "sgemm.h"
}

__global__ void fast_sgemm_gld_main_hw_ta_nb_fulltile (struct cublasSgemmParams parms)
{
#undef  FULL_TILES_ONLY
#undef  USE_MIXED_STEPPER
#undef  TRANSA
#undef  TRANSB
#undef  FAST_IMUL
#undef  USE_TEX
#define FULL_TILES_ONLY   1
#define USE_MIXED_STEPPER 0
#define TRANSA            1
#define TRANSB            0
#define FAST_IMUL         1
#define USE_TEX           0
#include "sgemm.h"
}

__global__ void fast_sgemm_gld_main_hw_ta_tb_fulltile (struct cublasSgemmParams parms)
{
#undef  FULL_TILES_ONLY
#undef  USE_MIXED_STEPPER
#undef  TRANSA
#undef  TRANSB
#undef  FAST_IMUL
#undef  USE_TEX
#define FULL_TILES_ONLY   1
#define USE_MIXED_STEPPER 0
#define TRANSA            1
#define TRANSB            1
#define FAST_IMUL         1
#define USE_TEX           0
#include "sgemm.h"
}

__global__ void fast_sgemm_gld_main_sw_na_nb (struct cublasSgemmParams parms)
{
#undef  FULL_TILES_ONLY
#undef  USE_MIXED_STEPPER
#undef  TRANSA
#undef  TRANSB
#undef  FAST_IMUL
#undef  USE_TEX
#define FULL_TILES_ONLY   0
#define USE_MIXED_STEPPER 1
#define TRANSA            0
#define TRANSB            0
#define FAST_IMUL         1
#define USE_TEX           0
#include "sgemm.h"
}

__global__ void fast_sgemm_gld_main_sw_na_tb (struct cublasSgemmParams parms)
{
#undef  FULL_TILES_ONLY
#undef  USE_MIXED_STEPPER
#undef  TRANSA
#undef  TRANSB
#undef  FAST_IMUL
#undef  USE_TEX
#define FULL_TILES_ONLY   0
#define USE_MIXED_STEPPER 1
#define TRANSA            0
#define TRANSB            1
#define FAST_IMUL         1
#define USE_TEX           0
#include "sgemm.h"
}

__global__ void fast_sgemm_gld_main_sw_ta_nb (struct cublasSgemmParams parms)
{
#undef  FULL_TILES_ONLY
#undef  USE_MIXED_STEPPER
#undef  TRANSA
#undef  TRANSB
#undef  FAST_IMUL
#undef  USE_TEX
#define FULL_TILES_ONLY   0
#define USE_MIXED_STEPPER 1
#define TRANSA            1
#define TRANSB            0
#define FAST_IMUL         1
#define USE_TEX           0
#include "sgemm.h"
}

__global__ void fast_sgemm_gld_main_sw_ta_tb (struct cublasSgemmParams parms)
{
#undef  FULL_TILES_ONLY
#undef  USE_MIXED_STEPPER
#undef  TRANSA
#undef  TRANSB
#undef  FAST_IMUL
#undef  USE_TEX
#define FULL_TILES_ONLY   0
#define USE_MIXED_STEPPER 1
#define TRANSA            1
#define TRANSB            1
#define FAST_IMUL         1
#define USE_TEX           0
#include "sgemm.h"
}

__global__ void fast_sgemm_gld_main_sw_na_nb_fulltile (struct cublasSgemmParams parms)
{
#undef  FULL_TILES_ONLY
#undef  USE_MIXED_STEPPER
#undef  TRANSA
#undef  TRANSB
#undef  FAST_IMUL
#undef  USE_TEX
#define FULL_TILES_ONLY   1
#define USE_MIXED_STEPPER 1
#define TRANSA            0
#define TRANSB            0
#define FAST_IMUL         1
#define USE_TEX           0
#include "sgemm.h"
}

__global__ void fast_sgemm_gld_main_sw_na_tb_fulltile (struct cublasSgemmParams parms)
{
#undef  FULL_TILES_ONLY
#undef  USE_MIXED_STEPPER
#undef  TRANSA
#undef  TRANSB
#undef  FAST_IMUL
#undef  USE_TEX
#define FULL_TILES_ONLY   1
#define USE_MIXED_STEPPER 1
#define TRANSA            0
#define TRANSB            1
#define FAST_IMUL         1
#define USE_TEX           0
#include "sgemm.h"
}

__global__ void fast_sgemm_gld_main_sw_ta_nb_fulltile (struct cublasSgemmParams parms)
{
#undef  FULL_TILES_ONLY
#undef  USE_MIXED_STEPPER
#undef  TRANSA
#undef  TRANSB
#undef  FAST_IMUL
#undef  USE_TEX
#define FULL_TILES_ONLY   1
#define USE_MIXED_STEPPER 1
#define TRANSA            1
#define TRANSB            0
#define FAST_IMUL         1
#define USE_TEX           0
#include "sgemm.h"
}

__global__ void fast_sgemm_gld_main_sw_ta_tb_fulltile (struct cublasSgemmParams parms)
{
#undef  FULL_TILES_ONLY
#undef  USE_MIXED_STEPPER
#undef  TRANSA
#undef  TRANSB
#undef  FAST_IMUL
#undef  USE_TEX
#define FULL_TILES_ONLY   1
#define USE_MIXED_STEPPER 1
#define TRANSA            1
#define TRANSB            1
#define FAST_IMUL         1
#define USE_TEX           0
#include "sgemm.h"
}