/*
__kernel void %PREFIXrotm_kernel( __global %TYPE *_X, __global %TYPE *_Y, uint N,
uint offx, int incx, uint offy, int incy
#ifndef DO_ROT
, __global %TYPE *_param, uint offParam // Rotm parameters
#else
, %PTYPE C, %PTYPE S // Rot parameters
#endif
*/
static void
assignKargs(KernelArg *args, const void *params, const void* )
{
CLBlasKargs *blasArgs = (CLBlasKargs*)params;
cl_int incx, incy;
INIT_KARG(&args[0], blasArgs->A);
INIT_KARG(&args[1], blasArgs->B);
initSizeKarg(&args[2], blasArgs->N);
initSizeKarg(&args[3], blasArgs->offBX);
incx = blasArgs->ldb.Vector;
INIT_KARG(&args[4], incx);
initSizeKarg(&args[5], blasArgs->offCY);
incy = blasArgs->ldc.Vector;
INIT_KARG(&args[6], incy);
if(blasArgs->pigFuncID == CLBLAS_ROT)
{
DataType alphaBetaType = (blasArgs->dtype == TYPE_COMPLEX_FLOAT)? TYPE_FLOAT:
((blasArgs->dtype == TYPE_COMPLEX_DOUBLE)? TYPE_DOUBLE: blasArgs->dtype);
assignScalarKarg(&args[7], &(blasArgs->alpha), alphaBetaType);
assignScalarKarg(&args[8], &(blasArgs->beta), alphaBetaType);
}
else if(blasArgs->pigFuncID == CLBLAS_ROTM)
{
INIT_KARG(&args[7], blasArgs->D);
initSizeKarg(&args[8], blasArgs->offd);
}
return;
}
static void
assignKargs(KernelArg *args, const void *params, const void*)
{
CLBlasKargs *blasArgs = (CLBlasKargs*)params;
#ifdef DEBUG_SYMM
printf("SAlpha=%f, DAlpha=%f, CAlpha =<%f, %f>, DAlpha=<%f, %f>\n",
blasArgs->alpha.argFloat, blasArgs->alpha.argDouble, CREAL(blasArgs->alpha.argFloatComplex), CIMAG(blasArgs->alpha.argFloatComplex),
CREAL(blasArgs->alpha.argDoubleComplex) , CIMAG(blasArgs->alpha.argDoubleComplex));
printf("SBeta=%f, DBeta=%f, CBeta=<%f, %f>, DBeta=<%f, %f>\n",
blasArgs->beta.argFloat, blasArgs->beta.argDouble, CREAL(blasArgs->beta.argFloatComplex), CIMAG(blasArgs->beta.argFloatComplex),
CREAL(blasArgs->beta.argDoubleComplex) , CIMAG(blasArgs->beta.argDoubleComplex));
#endif
INIT_KARG(&args[0], blasArgs->A); //A - input matrix - argument
INIT_KARG(&args[1], blasArgs->B);
INIT_KARG(&args[2], blasArgs->C);
initSizeKarg(&args[3], blasArgs->M);
initSizeKarg(&args[4], blasArgs->N);
initSizeKarg(&args[5], blasArgs->lda.matrix);
initSizeKarg(&args[6], blasArgs->ldb.matrix);
initSizeKarg(&args[7], blasArgs->ldc.matrix);
initSizeKarg(&args[8], blasArgs->offa); //PENDING: offA or offa ??
initSizeKarg(&args[9], blasArgs->offBX);
initSizeKarg(&args[10], blasArgs->offCY);
assignScalarKarg(&args[11], &(blasArgs->alpha), blasArgs->dtype);
assignScalarKarg(&args[12], &(blasArgs->beta), blasArgs->dtype);
return;
}
/*
__kernel void %PREFIXgbmv_RNT_kernel( __global const %TYPE * _A, __global %TYPE * _y_vector, __global %TYPE const* restrict _x_vector,
uint M, uint N, uint KL, uint KU, uint lda, int incx, int incy, uint offa, uint offx, uint offy
ifndef TBMV_ONLY
,%TYPE alpha, %TYPE beta
endif
*/
static void
assignKargs(KernelArg *args, const void *params, const void* )
{
CLBlasKargs *blasArgs = (CLBlasKargs*)params;
size_t fM, fN, fKL, fKU;
cl_int inc;
if( blasArgs->order == clblasColumnMajor ) // M, N, KL, KU gets swapped
{
fM = blasArgs->N;
fN = blasArgs->M;
fKL = blasArgs->KU;
fKU = blasArgs->KL;
}
else {
fM = blasArgs->M;
fN = blasArgs->N;
fKL = blasArgs->KL;
fKU = blasArgs->KU;
}
INIT_KARG(&args[0], blasArgs->A); //A - input matrix - argument
INIT_KARG(&args[1], blasArgs->C); //y - y vector
INIT_KARG(&args[2], blasArgs->B); //x - actual x vector argument
initSizeKarg(&args[3], fM);
initSizeKarg(&args[4], fN);
initSizeKarg(&args[5], fKL);
initSizeKarg(&args[6], fKU);
initSizeKarg(&args[7], blasArgs->lda.matrix);
inc = blasArgs->ldb.vector;
INIT_KARG(&args[8], inc);
inc = blasArgs->ldc.vector;
INIT_KARG(&args[9], inc);
initSizeKarg(&args[10], blasArgs->offa);
initSizeKarg(&args[11], blasArgs->offBX);
initSizeKarg(&args[12], blasArgs->offCY);
// For GBMV, SBMV, HBMV both alpha and beta has to be passed.
if( (blasArgs->pigFuncID == CLBLAS_GBMV) || (blasArgs->pigFuncID == CLBLAS_SBMV) || (blasArgs->pigFuncID == CLBLAS_HBMV) )
{
assignScalarKarg(&args[13], &(blasArgs->alpha), blasArgs->dtype);
assignScalarKarg(&args[14], &(blasArgs->beta), blasArgs->dtype);
}
#ifdef DEBUG_GBMV
printf("KL %d\tKU %d\n", fKL, fKU);
#endif
return;
}
/*
(__global %TYPE const* restrict A, __global %TYPE * _xnew, __global %TYPE const* restrict _x_vector, uint N,
int incx, int isUnity, uint lda, int doConj, uint offa, uint offx)
*/
static void
assignKargs(KernelArg *args, const void *params, const void* )
{
CLBlasKargs *blasArgs = (CLBlasKargs*)params;
//NOTE: This will not work if SolutionStep->args is not passed in const void *params.
SolutionStep *step = container_of(blasArgs, args, SolutionStep);
cl_int inc;
cl_int unity, doConj;
//bool incxOne = (blasArgs->ldb.vector == 1);
//bool incyOne = (blasArgs->ldc.vector == 1);
INIT_KARG(&args[0], blasArgs->A); //A - input matrix - argument
if( (step->funcID == CLBLAS_HEMV) || (blasArgs->pigFuncID == CLBLAS_HPMV) || (blasArgs->pigFuncID == CLBLAS_SPMV) )
{
INIT_KARG(&args[1], blasArgs->C); //y - since the 2nd argument is the result buffer, we should send y for HEMV
INIT_KARG(&args[2], blasArgs->B); //x - actual x vector argument
}
else
{
INIT_KARG(&args[1], blasArgs->B); //x - result buffer = _xnew argument
INIT_KARG(&args[2], blasArgs->C); //y - scratch == _x_vector argument
}
initSizeKarg(&args[3], blasArgs->N);
inc = blasArgs->ldb.vector;
INIT_KARG(&args[4], inc);
unity = (blasArgs->diag == clblasUnit);
INIT_KARG(&args[5], unity);
initSizeKarg(&args[6], blasArgs->lda.matrix);
doConj = (blasArgs->transA == clblasConjTrans);
#ifdef DEBUG_TRMV
printf("doConj is : %d, unity is : %d, incx is : %d\n", doConj, unity, inc);
#endif
INIT_KARG(&args[7], doConj);
initSizeKarg(&args[8], blasArgs->offa);
initSizeKarg(&args[9], blasArgs->offBX);
// For HEMV both alpha and beta has to be passed.
if( (step->funcID == CLBLAS_HEMV) || (blasArgs->pigFuncID == CLBLAS_HPMV) || (blasArgs->pigFuncID == CLBLAS_SPMV) )
{
inc = blasArgs->ldc.vector;
INIT_KARG(&args[10], inc);
initSizeKarg(&args[11], blasArgs->offCY);
assignScalarKarg(&args[12], &(blasArgs->alpha), blasArgs->dtype);
assignScalarKarg(&args[13], &(blasArgs->beta), blasArgs->dtype);
}
return;
}
static void
assignKargs(KernelArg *args, const void *params, const void *extra)
{
const CLBlasKargs *blasArgs = (const CLBlasKargs*)params;
KernelExtraFlags kflags = ((const CLBLASKernExtra*)extra)->flags;
int idx;
(void)extra;
initSizeKarg(&args[0], blasArgs->M);
initSizeKarg(&args[1], blasArgs->N);
assignScalarKarg(&args[2], &(blasArgs->alpha), blasArgs->dtype);
initMemobjKarg(&args[3], blasArgs->A, NULL, 0, 0);
initSizeKarg(&args[4], blasArgs->lda.matrix);
initMemobjKarg(&args[5], blasArgs->B, NULL, 0, 0);
initMemobjKarg(&args[6], blasArgs->B, NULL, 0, 0); //C in kernel
initSizeKarg(&args[7], blasArgs->ldb.matrix);
idx = 8;
if (kflags & KEXTRA_A_OFF_NOT_ZERO) {
initSizeKarg(&args[idx++], blasArgs->offA);
}
if (kflags & KEXTRA_BX_OFF_NOT_ZERO) {
initSizeKarg(&args[idx], blasArgs->offBX);
}
}
static void
assignKargs(KernelArg *args, const void *params, const void *extra)
{
const CLBlasKargs *blasArgs = (const CLBlasKargs*)params;
(void)extra;
switch (blasArgs->kernType) {
case CLBLAS_COMPUTING_KERNEL:
// arguments for computational kernel
initSizeKarg(&args[0], blasArgs->M);
initSizeKarg(&args[1], blasArgs->N);
initSizeKarg(&args[2], blasArgs->K);
assignScalarKarg(&args[3], &(blasArgs->alpha), blasArgs->dtype);
INIT_KARG(&args[4], blasArgs->scimage[0]);
INIT_KARG(&args[5], blasArgs->scimage[1]);
assignScalarKarg(&args[6], &(blasArgs->beta), blasArgs->dtype);
initMemobjKarg(&args[7], blasArgs->C, NULL, 0, 0);
initSizeKarg(&args[8], blasArgs->ldc.matrix);
initSizeKarg(&args[9], blasArgs->offCY);
break;
case CLBLAS_PREP_A_KERNEL:
INIT_KARG(&args[0], blasArgs->order);
INIT_KARG(&args[1], blasArgs->transA);
initSizeKarg(&args[2], blasArgs->M);
initSizeKarg(&args[3], blasArgs->K);
initMemobjKarg(&args[4], blasArgs->A, NULL, 0, 0);
initSizeKarg(&args[5], blasArgs->lda.matrix);
INIT_KARG(&args[6], blasArgs->scimage[0]);
initSizeKarg(&args[7], blasArgs->offA);
break;
case CLBLAS_PREP_B_KERNEL:
INIT_KARG(&args[0], blasArgs->order);
INIT_KARG(&args[1], blasArgs->transB);
initSizeKarg(&args[2], blasArgs->N);
initSizeKarg(&args[3], blasArgs->K);
initMemobjKarg(&args[4], blasArgs->B, NULL, 0, 0);
initSizeKarg(&args[5], blasArgs->ldb.matrix);
INIT_KARG(&args[6], blasArgs->scimage[1]);
initSizeKarg(&args[7], blasArgs->offBX);
break;
default:
//this should not happen
break;
}
}
static void
assignKargs(KernelArg *args, const void *params, const void *extra)
{
const CLBlasKargs *blasArgs = (const CLBlasKargs*)params;
KernelExtraFlags kflags = ((const CLBLASKernExtra*)extra)->flags;
cl_int inc;
int i;
initSizeKarg(&args[0], blasArgs->K);
assignScalarKarg(&args[1], &(blasArgs->alpha), blasArgs->dtype);
INIT_KARG(&args[2], blasArgs->A);
INIT_KARG(&args[3], blasArgs->B);
i = 4;
if (!(kflags & KEXTRA_BETA_ZERO)) {
assignScalarKarg(&args[i++], &(blasArgs->beta), blasArgs->dtype);
}
initMemobjKarg(&args[i++], blasArgs->C, NULL, 0, 0);
initSizeKarg(&args[i++], blasArgs->lda.matrix);
if (kflags & KEXTRA_A_OFF_NOT_ZERO) {
initSizeKarg(&args[i++], blasArgs->offA);
}
if (kflags & KEXTRA_BX_OFF_NOT_ZERO) {
initSizeKarg(&args[i++], blasArgs->offBX);
}
if (kflags & KEXTRA_CY_OFF_NOT_ZERO) {
initSizeKarg(&args[i++], blasArgs->offCY);
}
if (!(kflags & KEXTRA_INCX_ONE)) {
inc = blasArgs->ldb.vector;
INIT_KARG(&args[i], inc);
i++;
}
if (!(kflags & KEXTRA_INCY_ONE)) {
inc = blasArgs->ldc.vector;
INIT_KARG(&args[i], inc);
i++;
}
initSizeKarg(&args[i++], blasArgs->offsetN);
initSizeKarg(&args[i++], blasArgs->N); //Actual N
}
/*
( __global %TYPE* _A, __global const %TYPE* _X, int N,
int offx, int incx, int offa, int lda, %PTYPE alpha )
*/
static void
assignKargs(KernelArg *args, const void *params, const void*)
{
CLBlasKargs *blasArgs = (CLBlasKargs*)params;
cl_int incx;
INIT_KARG(&args[0], blasArgs->A); //A - input/output matrix - argument
INIT_KARG(&args[1], blasArgs->B); //x - x vector
initSizeKarg(&args[2], blasArgs->N);
initSizeKarg(&args[3], blasArgs->offBX);
incx = blasArgs->ldb.vector;
INIT_KARG(&args[4], incx);
initSizeKarg(&args[5], blasArgs->offa);
initSizeKarg(&args[6], blasArgs->lda.matrix);
DataType alphaType = (blasArgs->dtype == TYPE_COMPLEX_FLOAT)? TYPE_FLOAT : TYPE_DOUBLE;
assignScalarKarg(&args[7], &(blasArgs->alpha), alphaType);
return;
}
/*
__kernel void %PREFIXaxpy_kernel( %TYPE _alpha, __global %TYPE *_X, __global %TYPE *_Y,
uint N, uint offx, int incx, uint offy, int incy )
*/
static void
assignKargs(KernelArg *args, const void *params, const void* )
{
CLBlasKargs *blasArgs = (CLBlasKargs*)params;
cl_int incx, incy;
assignScalarKarg(&args[0], &(blasArgs->alpha), blasArgs->dtype);
INIT_KARG(&args[1], blasArgs->A);
INIT_KARG(&args[2], blasArgs->B);
initSizeKarg(&args[3], blasArgs->N);
initSizeKarg(&args[4], blasArgs->offBX);
incx = blasArgs->ldb.vector;
INIT_KARG(&args[5], incx);
initSizeKarg(&args[6], blasArgs->offCY);
incy = blasArgs->ldc.vector;
INIT_KARG(&args[7], incy);
return;
}
/*
( __global %TYPE* _A, __global const %TYPE* _X, __global const %TYPE* _Y, int N,
int offx, int incx, int offy, int incy, int offa, int lda, %TYPE alpha)
*/
static void
assignKargs(KernelArg *args, const void *params, const void*)
{
CLBlasKargs *blasArgs = (CLBlasKargs*)params;
cl_int inc;
INIT_KARG(&args[0], blasArgs->A); //A - input/output matrix - argument
INIT_KARG(&args[1], blasArgs->B); //X - x vector
INIT_KARG(&args[2], blasArgs->C); //Y - y vector
initSizeKarg(&args[3], blasArgs->N);
initSizeKarg(&args[4], blasArgs->offBX);
inc = blasArgs->ldb.vector;
INIT_KARG(&args[5], inc);
initSizeKarg(&args[6], blasArgs->offCY);
inc = blasArgs->ldc.vector;
INIT_KARG(&args[7], inc);
initSizeKarg(&args[8], blasArgs->offa);
initSizeKarg(&args[9], blasArgs->lda.matrix);
assignScalarKarg(&args[10], &(blasArgs->alpha), blasArgs->dtype);
return;
}
