Skip to content

Commit

Permalink
Merge pull request #22 from ecrc/acharara/cuda9
Browse files Browse the repository at this point in the history
fix compatibility with CUDA9, may need a patch for 'eps'
  • Loading branch information
acharara authored Oct 29, 2018
2 parents e38d63e + b5b26b6 commit 570a925
Show file tree
Hide file tree
Showing 47 changed files with 70 additions and 70 deletions.
2 changes: 1 addition & 1 deletion include/operators.h
Original file line number Diff line number Diff line change
Expand Up @@ -75,7 +75,7 @@ __device__ static __inline__ cuDoubleComplex make_real(cuDoubleComplex a){return

//==============================================================================================
#if defined(__CUDACC__)
#if (SM >= 30)
#if (TARGET_SM >= 30)
__device__ __inline__ float shfl(float x, int lane, int ws = 32)
{
return __shfl(x, lane, ws);
Expand Down
6 changes: 3 additions & 3 deletions make.inc
Original file line number Diff line number Diff line change
Expand Up @@ -46,7 +46,7 @@ LIB_KBLAS_NAME=kblas-gpu
NVOPTS =-O3 --compiler-options -fno-strict-aliasing
COPTS = -std=c99 -openmp

NVOPTS_3 = -DSM=$(_CUDA_ARCH_) -arch sm_$(_CUDA_ARCH_) -Xcompiler -fopenmp
NVOPTS_3 = -DTARGET_SM=$(_CUDA_ARCH_) -arch sm_$(_CUDA_ARCH_) -Xcompiler -fopenmp

#-----------------------------------------
ifdef _SUPPORT_SVD_
Expand All @@ -68,9 +68,9 @@ ifdef _SUPPORT_BLAS2_
GPU_ARCH_2=kepler

ifeq (${GPU_ARCH_2}, kepler)
NVOPTS_2 = -DSM=35 -arch sm_35
NVOPTS_2 = -DTARGET_SM=35 -arch sm_35
else ifeq (${GPU_ARCH_2}, fermi)
NVOPTS_2 = -DSM=20 -arch sm_20
NVOPTS_2 = -DTARGET_SM=20 -arch sm_20
else
$(error GPU_ARCH_2, currently ${GPU_TARGET_2}, must be either fermi or kepler for BLAS2 routines. Please edit your make.inc file)
endif
Expand Down
2 changes: 1 addition & 1 deletion src/batch_triangular/Xgemm_batch_core.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,7 @@
//==============================================================================================

//shuffle intrinsic is not supported before KEPLER
#if (SM >= 30)
#if (TARGET_SM >= 30)

//==============================================================================================
#define WARP 32
Expand Down
6 changes: 3 additions & 3 deletions src/batch_triangular/Xlauum_batch_kernels.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -28,9 +28,9 @@
//==============================================================================================
//Naming convention <dev/kernel>_<KernelName>_<Non/Uniform>_<Right/Left><Lower/Upper><Non/Transpose><Non/Diag>_<variants>
//==============================================================================================
#ifndef SM
#error "SM is not defined"
#elif (SM >= 30)
#ifndef TARGET_SM
#error "TARGET_SM is not defined"
#elif (TARGET_SM >= 30)


//==============================================================================================
Expand Down
6 changes: 3 additions & 3 deletions src/batch_triangular/Xposv_batch_kernels.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -28,9 +28,9 @@
//==============================================================================================
//Naming convention <dev/kernel>_<KernelName>_<Non/Uniform>_<Right/Left><Lower/Upper><Non/Transpose><Non/Diag>_<variants>
//==============================================================================================
#ifndef SM
#error "SM is not defined"
#elif (SM >= 30)
#ifndef TARGET_SM
#error "TARGET_SM is not defined"
#elif (TARGET_SM >= 30)

//==============================================================================================
template<typename T, int TX>
Expand Down
6 changes: 3 additions & 3 deletions src/batch_triangular/Xpoti_batch_kernels.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -28,9 +28,9 @@
//==============================================================================================
//Naming convention <dev/kernel>_<KernelName>_<Non/Uniform>_<Right/Left><Lower/Upper><Non/Transpose><Non/Diag>_<variants>
//==============================================================================================
#ifndef SM
#error "SM is not defined"
#elif (SM >= 30)
#ifndef TARGET_SM
#error "TARGET_SM is not defined"
#elif (TARGET_SM >= 30)

//==============================================================================================
template<typename T, int TX>
Expand Down
6 changes: 3 additions & 3 deletions src/batch_triangular/Xpotrf_batch_kernels.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -28,9 +28,9 @@
//==============================================================================================
//Naming convention <dev/kernel>_<KernelName>_<Non/Uniform>_<Right/Left><Lower/Upper><Non/Transpose><Non/Diag>_<variants>
//==============================================================================================
#ifndef SM
#error "SM is not defined"
#elif (SM >= 30)
#ifndef TARGET_SM
#error "TARGET_SM is not defined"
#elif (TARGET_SM >= 30)

//==============================================================================================

Expand Down
6 changes: 3 additions & 3 deletions src/batch_triangular/Xpotri_batch_kernels.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -28,9 +28,9 @@
//==============================================================================================
//Naming convention <dev/kernel>_<KernelName>_<Non/Uniform>_<Right/Left><Lower/Upper><Non/Transpose><Non/Diag>_<variants>
//==============================================================================================
#ifndef SM
#error "SM is not defined"
#elif (SM >= 30)
#ifndef TARGET_SM
#error "TARGET_SM is not defined"
#elif (TARGET_SM >= 30)

//==============================================================================================
template<typename T, int TX>
Expand Down
6 changes: 3 additions & 3 deletions src/batch_triangular/Xpotrs_batch_kernels.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -28,9 +28,9 @@
//==============================================================================================
//Naming convention <dev/kernel>_<KernelName>_<Non/Uniform>_<Right/Left><Lower/Upper><Non/Transpose><Non/Diag>_<variants>
//==============================================================================================
#ifndef SM
#error "SM is not defined"
#elif (SM >= 30)
#ifndef TARGET_SM
#error "TARGET_SM is not defined"
#elif (TARGET_SM >= 30)

//==============================================================================================
template<typename T, int TX>
Expand Down
6 changes: 3 additions & 3 deletions src/batch_triangular/Xsyrk_batch_kernels.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -28,12 +28,12 @@
//==============================================================================================
//Naming convention <dev/kernel>_<KernelName>_<Non/Uniform><Non/Strided>_<Lower/Upper><Non/Transpose>_<variants>
//==============================================================================================
#ifndef SM
#error "SM is not defined"
#ifndef TARGET_SM
#error "TARGET_SM is not defined"
#endif

//shuffle intrinsic is not supported before KEPLER
#if (SM >= 30)
#if (TARGET_SM >= 30)

//==============================================================================================
template<typename T, int B_ROWS, int A_COLS_PTY>
Expand Down
6 changes: 3 additions & 3 deletions src/batch_triangular/Xtrmm_batch_kernels.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -28,9 +28,9 @@
//==============================================================================================
//Naming convention <dev/kernel>_<KernelName>_<Non/Uniform>_<Right/Left><Lower/Upper><Non/Transpose><Non/Diag>_<variants>
//==============================================================================================
#ifndef SM
#error "SM is not defined"
#elif (SM >= 30)
#ifndef TARGET_SM
#error "TARGET_SM is not defined"
#elif (TARGET_SM >= 30)


//==============================================================================================
Expand Down
6 changes: 3 additions & 3 deletions src/batch_triangular/Xtrsm_batch_kernels.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -28,9 +28,9 @@
//==============================================================================================
//Naming convention <dev/kernel>_<KernelName>_<Non/Uniform>_<Right/Left><Lower/Upper><Non/Transpose><Non/Diag>_<variants>
//==============================================================================================
#ifndef SM
#error "SM is not defined"
#elif (SM >= 30)
#ifndef TARGET_SM
#error "TARGET_SM is not defined"
#elif (TARGET_SM >= 30)

//==============================================================================================
template<typename T, bool TRANS, int TX, int TY>
Expand Down
6 changes: 3 additions & 3 deletions src/batch_triangular/Xtrtri_batch_kernels.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -28,9 +28,9 @@
//==============================================================================================
//Naming convention <dev/kernel>_<KernelName>_<Non/Uniform>_<Right/Left><Lower/Upper><Non/Transpose><Non/Diag>_<variants>
//==============================================================================================
#ifndef SM
#error "SM is not defined"
#elif (SM >= 30)
#ifndef TARGET_SM
#error "TARGET_SM is not defined"
#elif (TARGET_SM >= 30)


//==============================================================================================
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/cgemv.cu
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,7 @@
#include <cublas.h>
#include "gemv_core.cuh"

#if (SM >= 30)
#if (TARGET_SM >= 30)

#define cgemvn_bs (32)
#define cgemvn_ty (16)
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/cgemv2.cu
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,7 @@
#include <cublas.h>
#include "gemv2_core.cuh"

#if(SM >= 30)
#if(TARGET_SM >= 30)

#define cgemvn_nb (32)
#define cgemvn_ntcol (4)
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/cgemv2_offset.cu
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,7 @@
#include <cublas.h>
#include "gemv2_offset_core.cuh"

#if(SM >= 30)
#if(TARGET_SM >= 30)

#define cgemvn_offset_nb (32)
#define cgemvn_offset_ntcol (2)
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/cgemv_mgpu.cu
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,7 @@
#include "defs.h"


#if(SM >= 30)
#if(TARGET_SM >= 30)

#define cgemvn_mgpu_bs (32)
#define cgemvn_mgpu_ty (4)
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/cgemv_offset.cu
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,7 @@
#include <cublas.h>
#include "gemv_offset_core.cuh"

#if(SM >= 30)
#if(TARGET_SM >= 30)

#define cgemvn_offset_bs (32)
#define cgemvn_offset_ty (8)
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/chemv.cu
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,7 @@

#include "syhemv_core.cuh"

#if(SM >= 30)
#if(TARGET_SM >= 30)

#define chemv_upper_bs (32)
#define chemv_upper_ty (2)
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/chemv_mgpu.cu
Original file line number Diff line number Diff line change
Expand Up @@ -20,7 +20,7 @@
#include "syhemv_mgpu_offset_core.cuh"
#include "defs.h"

#if(SM >= 30)
#if(TARGET_SM >= 30)

#define chemv_upper_bs (32)
#define chemv_upper_ty (4)
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/chemv_offset.cu
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,7 @@

#include "syhemv_offset_core.cuh"

#if(SM >= 30)
#if(TARGET_SM >= 30)

#define chemv_upper_bs (32)
#define chemv_upper_ty (2)
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/dgemv.cu
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,7 @@
#include <cublas.h>
#include "gemv_core.cuh"

#if(SM >= 30)
#if(TARGET_SM >= 30)

#define dgemvn_bs (64)
#define dgemvn_ty (4)
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/dgemv2.cu
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,7 @@
#include <cublas.h>
#include "gemv2_core.cuh"

#if(SM >= 30)
#if(TARGET_SM >= 30)

#define dgemvn_nb (32)
#define dgemvn_ntcol (4)
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/dgemv2_offset.cu
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,7 @@
#include <cublas.h>
#include "gemv2_offset_core.cuh"

#if(SM >= 30)
#if(TARGET_SM >= 30)

#define dgemvn_offset_nb (32)
#define dgemvn_offset_ntcol (4)
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/dgemv_mgpu.cu
Original file line number Diff line number Diff line change
Expand Up @@ -24,7 +24,7 @@
#include "gemv_mgpu_offset_core.cuh"
#include "defs.h"

#if(SM >= 30)
#if(TARGET_SM >= 30)

#define dgemvn_mgpu_bs (64)
#define dgemvn_mgpu_ty (4)
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/dgemv_offset.cu
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,7 @@
#include <cublas.h>
#include "gemv_offset_core.cuh"

#if(SM >= 30)
#if(TARGET_SM >= 30)

#define dgemvn_offset_bs (32)
#define dgemvn_offset_ty (8)
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/dsymv.cu
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,7 @@

#include "syhemv_core.cuh"

#if(SM >= 30)
#if(TARGET_SM >= 30)

#define dsymv_upper_bs (32)
#define dsymv_upper_ty (2)
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/dsymv_mgpu.cu
Original file line number Diff line number Diff line change
Expand Up @@ -20,7 +20,7 @@
#include "syhemv_mgpu_offset_core.cuh"
#include "defs.h"

#if(SM >= 30)
#if(TARGET_SM >= 30)

#define dsymv_upper_bs (32)
#define dsymv_upper_ty (4)
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/dsymv_offset.cu
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,7 @@

#include "syhemv_offset_core.cuh"

#if(SM >= 30)
#if(TARGET_SM >= 30)

#define dsymv_upper_bs (32)
#define dsymv_upper_ty (2)
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/sgemv.cu
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,7 @@
#include <cublas.h>
#include "gemv_core.cuh"

#if(SM >= 30)
#if(TARGET_SM >= 30)

#define sgemvn_bs (64)
#define sgemvn_ty (8)
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/sgemv2.cu
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,7 @@
#include <cublas.h>
#include "gemv2_core.cuh"

#if(SM >= 30)
#if(TARGET_SM >= 30)

#define sgemvn_nb (32)
#define sgemvn_ntcol (4)
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/sgemv2_offset.cu
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,7 @@
#include <cublas.h>
#include "gemv2_offset_core.cuh"

#if(SM >= 30)
#if(TARGET_SM >= 30)

#define sgemvn_offset_nb (64)
#define sgemvn_offset_ntcol (4)
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/sgemv_mgpu.cu
Original file line number Diff line number Diff line change
Expand Up @@ -24,7 +24,7 @@
#include "gemv_mgpu_offset_core.cuh"
#include "defs.h"

#if(SM >= 30)
#if(TARGET_SM >= 30)

#define sgemvn_mgpu_bs (64)
#define sgemvn_mgpu_ty (4)
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/sgemv_offset.cu
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,7 @@
#include <cublas.h>
#include "gemv_offset_core.cuh"

#if(SM >= 30)
#if(TARGET_SM >= 30)

#define sgemvn_offset_bs (64)
#define sgemvn_offset_ty (16)
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/ssymv.cu
Original file line number Diff line number Diff line change
Expand Up @@ -105,7 +105,7 @@

#include "syhemv_core.cuh"

#if(SM >= 30)
#if(TARGET_SM >= 30)

#define ssymv_upper_bs (64)
#define ssymv_upper_ty (4)
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/ssymv_mgpu.cu
Original file line number Diff line number Diff line change
Expand Up @@ -20,7 +20,7 @@
#include "syhemv_mgpu_offset_core.cuh"
#include "defs.h"

#if(SM >= 30)
#if(TARGET_SM >= 30)

#define ssymv_upper_bs (64)
#define ssymv_upper_ty (4)
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/ssymv_offset.cu
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,7 @@

#include "syhemv_offset_core.cuh"

#if(SM >= 30)
#if(TARGET_SM >= 30)

#define ssymv_upper_bs (64)
#define ssymv_upper_ty (4)
Expand Down
2 changes: 1 addition & 1 deletion src/blas_l2/zgemv.cu
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,7 @@
#include <cublas.h>
#include "gemv_core.cuh"

#if(SM >= 30)
#if(TARGET_SM >= 30)

#define zgemvn_bs (16)
#define zgemvn_ty (8)
Expand Down
Loading

0 comments on commit 570a925

Please sign in to comment.