static void create_task_22(starpu_data_handle_t dataA, unsigned k, unsigned i, unsigned j)
{
/* FPRINTF(stdout, "task 22 k,i,j = %d,%d,%d TAG = %llx\n", k,i,j, TAG22(k,i,j)); */
struct starpu_task *task = create_task(TAG22(k, i, j));
task->cl = &cl22;
/* which sub-data is manipulated ? */
task->handles[0] = starpu_data_get_sub_data(dataA, 2, k, i);
task->handles[1] = starpu_data_get_sub_data(dataA, 2, k, j);
task->handles[2] = starpu_data_get_sub_data(dataA, 2, i, j);
if (!noprio && (i == k + 1) && (j == k +1) )
{
task->priority = STARPU_MAX_PRIO;
}
/* enforce dependencies ... */
if (k > 0)
{
starpu_tag_declare_deps(TAG22(k, i, j), 3, TAG22(k-1, i, j), TAG21(k, i), TAG21(k, j));
}
else
{
starpu_tag_declare_deps(TAG22(k, i, j), 2, TAG21(k, i), TAG21(k, j));
}
int n = starpu_matrix_get_nx(task->handles[0]);
task->flops = FLOPS_SGEMM(n, n, n);
int ret = starpu_task_submit(task);
if (STARPU_UNLIKELY(ret == -ENODEV))
{
FPRINTF(stderr, "No worker may execute this task\n");
exit(0);
}
}
static int create_task_22(starpu_data_handle_t dataA, unsigned k, unsigned i, unsigned j, unsigned reclevel)
{
int ret;
/* FPRINTF(stdout, "task 22 k,i,j = %d,%d,%d TAG = %llx\n", k,i,j, TAG22_AUX(k,i,j)); */
struct starpu_task *task = create_task(TAG22_AUX(k, i, j, reclevel));
task->cl = &cl22;
/* which sub-data is manipulated ? */
task->handles[0] = starpu_data_get_sub_data(dataA, 2, k, i);
task->handles[1] = starpu_data_get_sub_data(dataA, 2, k, j);
task->handles[2] = starpu_data_get_sub_data(dataA, 2, i, j);
if ( (i == k + 1) && (j == k +1) )
{
task->priority = STARPU_MAX_PRIO;
}
/* enforce dependencies ... */
if (k > 0)
{
starpu_tag_declare_deps(TAG22_AUX(k, i, j, reclevel), 3, TAG22_AUX(k-1, i, j, reclevel), TAG21_AUX(k, i, reclevel), TAG21_AUX(k, j, reclevel));
}
else
{
starpu_tag_declare_deps(TAG22_AUX(k, i, j, reclevel), 2, TAG21_AUX(k, i, reclevel), TAG21_AUX(k, j, reclevel));
}
int n = starpu_matrix_get_nx(task->handles[0]);
task->flops = FLOPS_SGEMM(n, n, n);
ret = starpu_task_submit(task);
if (ret != -ENODEV) STARPU_CHECK_RETURN_VALUE(ret, "starpu_task_submit");
return ret;
}
int main( int argc, char** argv)
{
real_Double_t gflops, magma_perf, magma_time, clblas_perf, clblas_time, cpu_perf, cpu_time;
float magma_error, clblas_error, work[1];
magma_trans_t transA = MagmaNoTrans;
magma_trans_t transB = MagmaNoTrans;
magma_int_t istart = 1024;
magma_int_t iend = 6240;
magma_int_t M, M0 = 0;
magma_int_t N, N0 = 0;
magma_int_t K, K0 = 0;
magma_int_t i;
magma_int_t Am, An, Bm, Bn;
magma_int_t szeA, szeB, szeC;
magma_int_t lda, ldb, ldc, ldda, lddb, lddc;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
float *h_A, *h_B, *h_C, *h_C2, *h_C3;
magmaFloat_ptr d_A, d_B, d_C;
float c_neg_one = MAGMA_S_NEG_ONE;
float alpha = MAGMA_S_MAKE( 0.29, -0.86 );
float beta = MAGMA_S_MAKE( -0.48, 0.38 );
int lapack = getenv("MAGMA_RUN_LAPACK") != NULL;
int count = 1;
printf("\nUsage: testing_sgemm [-NN|NT|TN|TT|NC|CN|TC|CT|CC] -M m -N n -K k -count c -l\n"
" -l or setting $MAGMA_RUN_LAPACK runs CPU BLAS,\n"
" and computes both MAGMA and CLBLAS error using CPU BLAS result.\n"
" Else, MAGMA error is computed using CLBLAS result.\n\n");
for( int i = 1; i < argc; ++i ) {
if ( strcmp("-N", argv[i]) == 0 && i+1 < argc ){
N0 = atoi(argv[++i]);
}
else if ( strcmp("-M", argv[i]) == 0 && i+1 < argc ){
M0 = atoi(argv[++i]);
}
else if ( strcmp("-K", argv[i]) == 0 && i+1 < argc ){
K0 = atoi(argv[++i]);
}
else if (strcmp("-NN", argv[i])==0){
transA = transB = MagmaNoTrans;
}
else if (strcmp("-TT", argv[i])==0){
transA = transB = MagmaTrans;
}
else if (strcmp("-NT", argv[i])==0){
transA = MagmaNoTrans;
transB = MagmaTrans;
}
else if (strcmp("-TN", argv[i])==0){
transA = MagmaTrans;
transB = MagmaNoTrans;
}
else if (strcmp("-NC", argv[i])==0){
transA = MagmaNoTrans;
transB = MagmaConjTrans;
}
else if (strcmp("-TC", argv[i])==0){
transA = MagmaTrans;
transB = MagmaConjTrans;
}
else if (strcmp("-CN", argv[i])==0){
transA = MagmaConjTrans;
transB = MagmaNoTrans;
}
else if (strcmp("-CT", argv[i])==0){
transA = MagmaConjTrans;
transB = MagmaTrans;
}
else if (strcmp("-CC", argv[i])==0){
transA = transB = MagmaConjTrans;
}
else if (strcmp("-l", argv[i])==0) {
lapack = true;
}
else if ( strcmp("-count", argv[i]) == 0 && i+1 < argc ){
count = atoi(argv[++i]);
}
else {
printf( "invalid argument: %s\n", argv[i] );
exit(1);
}
}
if ( (M0 != 0) && (N0 != 0) && (K0 != 0) )
iend = istart + 1;
M = N = K = iend;
if ( M0 != 0 ) M = M0;
if ( N0 != 0 ) N = N0;
if ( K0 != 0 ) K = K0;
if( transA == MagmaNoTrans ) {
Am = M;
An = K;
} else {
Am = K;
An = M;
}
if( transB == MagmaNoTrans ) {
Bm = K;
Bn = N;
} else {
Bm = N;
Bn = K;
}
/* Initialize */
magma_queue_t queue;
magma_device_t device[ MagmaMaxGPUs ];
magma_int_t num = 0;
magma_int_t err;
magma_init();
err = magma_getdevices( device, MagmaMaxGPUs, &num );
if ( err != 0 || num < 1 ) {
fprintf( stderr, "magma_getdevices failed: %d\n", (int) err );
exit(-1);
}
err = magma_queue_create( device[0], &queue );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", (int) err );
exit(-1);
}
lda = ldc = M;
ldb = Bm;
ldda = ((M+31)/32)*32;
lddb = ((ldb+31)/32)*32;
lddc = ldda;
K += 32;
M += 32;
N += 32;
TESTING_MALLOC_CPU( h_A, float, lda*K );
TESTING_MALLOC_CPU( h_B, float, ldb*Bn );
TESTING_MALLOC_CPU( h_C, float, ldc*N );
TESTING_MALLOC_CPU( h_C2, float, ldc*N );
TESTING_MALLOC_CPU( h_C3, float, ldc*N );
TESTING_MALLOC_DEV( d_A, float, ldda*K );
TESTING_MALLOC_DEV( d_B, float, lddb*Bn );
TESTING_MALLOC_DEV( d_C, float, lddc*N );
printf("Testing transA = %c transB = %c\n", *lapack_const(transA), *lapack_const(transB));
printf(" M N K MAGMA Gflop/s (sec) CLBLAS Gflop/s (sec) CPU Gflop/s (sec) MAGMA error CLBLAS error\n");
printf("===========================================================================================================\n");
for( i=istart; i<iend; i = (int)(i*1.25) ) {
for( int cnt = 0; cnt < count; ++cnt ) {
M = N = K = i;
if ( M0 != 0 ) M = M0;
if ( N0 != 0 ) N = N0;
if ( K0 != 0 ) K = K0;
if( transA == MagmaNoTrans ) {
lda = Am = M;
An = K;
} else {
lda = Am = K;
An = M;
}
if( transB == MagmaNoTrans ) {
ldb = Bm = K;
Bn = N;
} else {
ldb = Bm = N;
Bn = K;
}
gflops = FLOPS_SGEMM( M, N, K ) / 1e9;
ldc = M;
ldda = ((lda+31)/32)*32;
lddb = ((ldb+31)/32)*32;
lddc = ((ldc+31)/32)*32;
szeA = lda * An;
szeB = ldb * Bn;
szeC = ldc * N;
/* Initialize the matrices */
lapackf77_slarnv( &ione, ISEED, &szeA, h_A );
lapackf77_slarnv( &ione, ISEED, &szeB, h_B );
lapackf77_slarnv( &ione, ISEED, &szeC, h_C );
/* =====================================================================
Performs operation using MAGMA-BLAS
=================================================================== */
magma_ssetmatrix( Am, An, h_A, lda, d_A, 0, ldda, queue );
magma_ssetmatrix( Bm, Bn, h_B, ldb, d_B, 0, lddb, queue );
magma_ssetmatrix( M, N, h_C, ldc, d_C, 0, lddc, queue );
magmablas_sgemm_reduce( M, N, K,
alpha, d_A, 0, ldda,
d_B, 0, lddb,
beta, d_C, 0, lddc, queue );
magma_ssetmatrix( M, N, h_C, ldc, d_C, 0, lddc, queue );
magma_queue_sync(queue);
magma_time = magma_wtime();
magmablas_sgemm_reduce( M, N, K,
alpha, d_A, 0, ldda,
d_B, 0, lddb,
beta, d_C, 0, lddc, queue );
magma_queue_sync(queue);
magma_time = magma_wtime() - magma_time;
magma_perf = gflops / magma_time;
magma_sgetmatrix( M, N, d_C, 0, lddc, h_C2, ldc, queue );
/* =====================================================================
Performs operation using CUDA-BLAS
=================================================================== */
magma_ssetmatrix( M, N, h_C, ldc, d_C, 0, lddc, queue );
magma_sgemm( transA, transB, M, N, K,
alpha, d_A, 0, ldda,
d_B, 0, lddb,
beta, d_C, 0, lddc, queue );
magma_ssetmatrix( M, N, h_C, ldc, d_C, 0, lddc, queue );
magma_queue_sync(queue);
clblas_time = magma_wtime();
magma_sgemm( transA, transB, M, N, K,
alpha, d_A, 0, ldda,
d_B, 0, lddb,
beta, d_C, 0, lddc, queue );
magma_queue_sync(queue);
clblas_time = magma_wtime() - clblas_time;
clblas_perf = gflops / clblas_time;
magma_sgetmatrix( M, N, d_C, 0, lddc, h_C3, ldc, queue );
/* =====================================================================
Performs operation using BLAS
=================================================================== */
if ( lapack ) {
cpu_time = magma_wtime();
blasf77_sgemm( lapack_const(transA), lapack_const(transB), &M, &N, &K,
&alpha, h_A, &lda,
h_B, &ldb,
&beta, h_C, &ldc );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
}
/* =====================================================================
Error Computation and Performance Compariosn
=================================================================== */
if ( lapack ) {
// compare both magma & clblas to lapack
blasf77_saxpy(&szeC, &c_neg_one, h_C, &ione, h_C2, &ione);
magma_error = lapackf77_slange("M", &M, &N, h_C2, &ldc, work);
blasf77_saxpy(&szeC, &c_neg_one, h_C, &ione, h_C3, &ione);
clblas_error = lapackf77_slange("M", &M, &N, h_C3, &ldc, work);
printf("%5d %5d %5d %7.2f (%7.4f) %7.2f (%7.4f) %7.2f (%7.4f) %8.2e %8.2e\n",
(int) M, (int) N, (int) K,
magma_perf, magma_time, clblas_perf, clblas_time, cpu_perf, cpu_time,
magma_error, clblas_error );
}
else {
// compare magma to clblas
blasf77_saxpy(&szeC, &c_neg_one, h_C3, &ione, h_C2, &ione);
magma_error = lapackf77_slange("M", &M, &N, h_C2, &ldc, work);
printf("%5d %5d %5d %7.2f (%7.4f) %7.2f (%7.4f) --- ( --- ) %8.2e ---\n",
(int) M, (int) N, (int) K,
magma_perf, magma_time, clblas_perf, clblas_time,
magma_error );
}
}
if ( count > 1 ) {
printf( "\n" );
}
}
/* Memory clean up */
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_C );
TESTING_FREE_CPU( h_C2 );
TESTING_FREE_CPU( h_C3 );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
TESTING_FREE_DEV( d_C );
magma_queue_destroy( queue );
magma_finalize();
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing sgemm
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, dev_perf, dev_time, cpu_perf, cpu_time;
float magma_error, dev_error, Cnorm, work[1];
magma_int_t M, N, K;
magma_int_t Am, An, Bm, Bn;
magma_int_t sizeA, sizeB, sizeC;
magma_int_t lda, ldb, ldc, ldda, lddb, lddc;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
float *h_A, *h_B, *h_C, *h_Cmagma, *h_Cdev;
magmaFloat_ptr d_A, d_B, d_C;
float c_neg_one = MAGMA_S_NEG_ONE;
float alpha = MAGMA_S_MAKE( 0.29, -0.86 );
float beta = MAGMA_S_MAKE( -0.48, 0.38 );
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
#ifdef HAVE_CUBLAS
// for CUDA, we can check MAGMA vs. CUBLAS, without running LAPACK
printf("If running lapack (option --lapack), MAGMA and %s error are both computed\n"
"relative to CPU BLAS result. Else, MAGMA error is computed relative to %s result.\n\n",
g_platform_str, g_platform_str );
printf("transA = %s, transB = %s\n",
lapack_trans_const(opts.transA),
lapack_trans_const(opts.transB) );
printf(" M N K MAGMA Gflop/s (ms) %s Gflop/s (ms) CPU Gflop/s (ms) MAGMA error %s error\n",
g_platform_str, g_platform_str );
#else
// for others, we need LAPACK for check
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
printf("transA = %s, transB = %s\n",
lapack_trans_const(opts.transA),
lapack_trans_const(opts.transB) );
printf(" M N K %s Gflop/s (ms) CPU Gflop/s (ms) %s error\n",
g_platform_str, g_platform_str );
#endif
printf("=========================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
K = opts.ksize[itest];
gflops = FLOPS_SGEMM( M, N, K ) / 1e9;
if ( opts.transA == MagmaNoTrans ) {
lda = Am = M;
An = K;
} else {
lda = Am = K;
An = M;
}
if ( opts.transB == MagmaNoTrans ) {
ldb = Bm = K;
Bn = N;
} else {
ldb = Bm = N;
Bn = K;
}
ldc = M;
ldda = ((lda+31)/32)*32;
lddb = ((ldb+31)/32)*32;
lddc = ((ldc+31)/32)*32;
sizeA = lda*An;
sizeB = ldb*Bn;
sizeC = ldc*N;
TESTING_MALLOC_CPU( h_A, float, lda*An );
TESTING_MALLOC_CPU( h_B, float, ldb*Bn );
TESTING_MALLOC_CPU( h_C, float, ldc*N );
TESTING_MALLOC_CPU( h_Cmagma, float, ldc*N );
TESTING_MALLOC_CPU( h_Cdev, float, ldc*N );
TESTING_MALLOC_DEV( d_A, float, ldda*An );
TESTING_MALLOC_DEV( d_B, float, lddb*Bn );
TESTING_MALLOC_DEV( d_C, float, lddc*N );
/* Initialize the matrices */
lapackf77_slarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_slarnv( &ione, ISEED, &sizeB, h_B );
lapackf77_slarnv( &ione, ISEED, &sizeC, h_C );
magma_ssetmatrix( Am, An, h_A, lda, d_A, 0, ldda, opts.queue );
magma_ssetmatrix( Bm, Bn, h_B, ldb, d_B, 0, lddb, opts.queue );
/* =====================================================================
Performs operation using MAGMABLAS (currently only with CUDA)
=================================================================== */
#ifdef HAVE_CUBLAS
magma_ssetmatrix( M, N, h_C, ldc, d_C, lddc );
magma_time = magma_sync_wtime( NULL );
magmablas_sgemm( opts.transA, opts.transB, M, N, K,
alpha, d_A, ldda,
d_B, lddb,
beta, d_C, lddc );
magma_time = magma_sync_wtime( NULL ) - magma_time;
magma_perf = gflops / magma_time;
magma_sgetmatrix( M, N, d_C, lddc, h_Cmagma, ldc );
#endif
/* =====================================================================
Performs operation using CUBLAS / clBLAS / Xeon Phi MKL
=================================================================== */
magma_ssetmatrix( M, N, h_C, ldc, d_C, 0, lddc, opts.queue );
#ifdef HAVE_CUBLAS
dev_time = magma_sync_wtime( NULL );
cublasSgemm( opts.handle, cublas_trans_const(opts.transA), cublas_trans_const(opts.transB), M, N, K,
&alpha, d_A, ldda,
d_B, lddb,
&beta, d_C, lddc );
dev_time = magma_sync_wtime( NULL ) - dev_time;
#else
dev_time = magma_sync_wtime( opts.queue );
magma_sgemm( opts.transA, opts.transB, M, N, K,
alpha, d_A, 0, ldda,
d_B, 0, lddb,
beta, d_C, 0, lddc, opts.queue );
dev_time = magma_sync_wtime( opts.queue ) - dev_time;
#endif
dev_perf = gflops / dev_time;
magma_sgetmatrix( M, N, d_C, 0, lddc, h_Cdev, ldc, opts.queue );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
blasf77_sgemm( lapack_trans_const(opts.transA), lapack_trans_const(opts.transB), &M, &N, &K,
&alpha, h_A, &lda,
h_B, &ldb,
&beta, h_C, &ldc );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
}
/* =====================================================================
Check the result
=================================================================== */
if ( opts.lapack ) {
// compute relative error for both magma & dev, relative to lapack,
// |C_magma - C_lapack| / |C_lapack|
Cnorm = lapackf77_slange( "F", &M, &N, h_C, &ldc, work );
blasf77_saxpy( &sizeC, &c_neg_one, h_C, &ione, h_Cdev, &ione );
dev_error = lapackf77_slange( "F", &M, &N, h_Cdev, &ldc, work ) / Cnorm;
#ifdef HAVE_CUBLAS
blasf77_saxpy( &sizeC, &c_neg_one, h_C, &ione, h_Cmagma, &ione );
magma_error = lapackf77_slange( "F", &M, &N, h_Cmagma, &ldc, work ) / Cnorm;
printf("%5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %s\n",
(int) M, (int) N, (int) K,
magma_perf, 1000.*magma_time,
dev_perf, 1000.*dev_time,
cpu_perf, 1000.*cpu_time,
magma_error, dev_error,
(magma_error < tol && dev_error < tol ? "ok" : "failed"));
status += ! (magma_error < tol && dev_error < tol);
#else
printf("%5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) M, (int) N, (int) K,
dev_perf, 1000.*dev_time,
cpu_perf, 1000.*cpu_time,
dev_error,
(dev_error < tol ? "ok" : "failed"));
status += ! (dev_error < tol);
#endif
}
else {
#ifdef HAVE_CUBLAS
// compute relative error for magma, relative to dev (currently only with CUDA)
Cnorm = lapackf77_slange( "F", &M, &N, h_Cdev, &ldc, work );
blasf77_saxpy( &sizeC, &c_neg_one, h_Cdev, &ione, h_Cmagma, &ione );
magma_error = lapackf77_slange( "F", &M, &N, h_Cmagma, &ldc, work ) / Cnorm;
printf("%5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) --- ( --- ) %8.2e --- %s\n",
(int) M, (int) N, (int) K,
magma_perf, 1000.*magma_time,
dev_perf, 1000.*dev_time,
magma_error,
(magma_error < tol ? "ok" : "failed"));
status += ! (magma_error < tol);
#else
printf("%5d %5d %5d %7.2f (%7.2f) --- ( --- ) ---\n",
(int) M, (int) N, (int) K,
dev_perf, 1000.*dev_time );
#endif
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_C );
TESTING_FREE_CPU( h_Cmagma );
TESTING_FREE_CPU( h_Cdev );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
TESTING_FREE_DEV( d_C );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, t1, t2;
float c_neg_one = MAGMA_S_NEG_ONE;
magma_int_t ione = 1;
const char trans[] = { 'N', 'C', 'T' };
const char uplo[] = { 'L', 'U' };
const char diag[] = { 'U', 'N' };
const char side[] = { 'L', 'R' };
float *A, *B, *C, *C2, *LU;
float *dA, *dB, *dC1, *dC2;
float alpha = MAGMA_S_MAKE( 0.5, 0.1 );
float beta = MAGMA_S_MAKE( 0.7, 0.2 );
float dalpha = 0.6;
float dbeta = 0.8;
float work[1], error, total_error;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t m, n, k, size, maxn, ld, info;
magma_int_t *piv;
magma_err_t err;
magma_opts opts;
parse_opts( argc, argv, &opts );
printf( "Compares magma wrapper function to cublas function; all diffs should be exactly 0.\n\n" );
total_error = 0.;
for( int i = 0; i < opts.ntest; ++i ) {
m = opts.msize[i];
n = opts.nsize[i];
k = opts.ksize[i];
printf("=========================================================================\n");
printf( "M %d, N %d, K %d\n", (int) m, (int) n, (int) k );
// allocate matrices
// over-allocate so they can be any combination of {m,n,k} x {m,n,k}.
maxn = max( max( m, n ), k );
ld = maxn;
size = maxn*maxn;
err = magma_malloc_cpu( (void**) &piv, maxn*sizeof(magma_int_t) ); assert( err == 0 );
err = magma_smalloc_pinned( &A, size ); assert( err == 0 );
err = magma_smalloc_pinned( &B, size ); assert( err == 0 );
err = magma_smalloc_pinned( &C, size ); assert( err == 0 );
err = magma_smalloc_pinned( &C2, size ); assert( err == 0 );
err = magma_smalloc_pinned( &LU, size ); assert( err == 0 );
err = magma_smalloc( &dA, size ); assert( err == 0 );
err = magma_smalloc( &dB, size ); assert( err == 0 );
err = magma_smalloc( &dC1, size ); assert( err == 0 );
err = magma_smalloc( &dC2, size ); assert( err == 0 );
// initialize matrices
size = maxn*maxn;
lapackf77_slarnv( &ione, ISEED, &size, A );
lapackf77_slarnv( &ione, ISEED, &size, B );
lapackf77_slarnv( &ione, ISEED, &size, C );
printf( "========== Level 1 BLAS ==========\n" );
// ----- test SSWAP
// swap 2nd and 3rd columns of dA, then copy to C2 and compare with A
assert( n >= 4 );
magma_ssetmatrix( m, n, A, ld, dA, ld );
magma_ssetmatrix( m, n, A, ld, dB, ld );
magma_sswap( m, dA(0,1), 1, dA(0,2), 1 );
magma_sswap( m, dB(0,1), 1, dB(0,2), 1 );
// check results, storing diff between magma and cuda calls in C2
cublasSaxpy( ld*n, c_neg_one, dA, 1, dB, 1 );
magma_sgetmatrix( m, n, dB, ld, C2, ld );
error = lapackf77_slange( "F", &m, &k, C2, &ld, work );
total_error += error;
printf( "sswap diff %.2g\n", error );
// ----- test ISAMAX
// get argmax of column of A
magma_ssetmatrix( m, k, A, ld, dA, ld );
error = 0;
for( int j = 0; j < k; ++j ) {
magma_int_t i1 = magma_isamax( m, dA(0,j), 1 );
magma_int_t i2 = cublasIsamax( m, dA(0,j), 1 );
assert( i1 == i2 );
error += abs( i1 - i2 );
}
total_error += error;
gflops = (float)m * k / 1e9;
printf( "isamax diff %.2g\n", error );
printf( "\n" );
printf( "========== Level 2 BLAS ==========\n" );
// ----- test SGEMV
// c = alpha*A*b + beta*c, with A m*n; b,c m or n-vectors
// try no-trans/trans
for( int ia = 0; ia < 3; ++ia ) {
magma_ssetmatrix( m, n, A, ld, dA, ld );
magma_ssetvector( maxn, B, 1, dB, 1 );
magma_ssetvector( maxn, C, 1, dC1, 1 );
magma_ssetvector( maxn, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_sgemv( trans[ia], m, n, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasSgemv( trans[ia], m, n, alpha, dA, ld, dB, 1, beta, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
size = (trans[ia] == 'N' ? m : n);
cublasSaxpy( size, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetvector( size, dC2, 1, C2, 1 );
error = lapackf77_slange( "F", &size, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_SGEMV( m, n ) / 1e9;
printf( "sgemv( %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
trans[ia], error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test SSYMV
// c = alpha*A*b + beta*c, with A m*m symmetric; b,c m-vectors
// try upper/lower
for( int iu = 0; iu < 2; ++iu ) {
magma_ssetmatrix( m, m, A, ld, dA, ld );
magma_ssetvector( m, B, 1, dB, 1 );
magma_ssetvector( m, C, 1, dC1, 1 );
magma_ssetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_ssymv( uplo[iu], m, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasSsymv( uplo[iu], m, alpha, dA, ld, dB, 1, beta, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasSaxpy( m, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_slange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_SSYMV( m ) / 1e9;
printf( "ssymv( %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test STRSV
// solve A*c = c, with A m*m triangular; c m-vector
// try upper/lower, no-trans/trans, unit/non-unit diag
// Factor A into LU to get well-conditioned triangles, else solve yields garbage.
// Still can give garbage if solves aren't consistent with LU factors,
// e.g., using unit diag for U, so copy lower triangle to upper triangle.
// Also used for trsm later.
lapackf77_slacpy( "Full", &maxn, &maxn, A, &ld, LU, &ld );
lapackf77_sgetrf( &maxn, &maxn, LU, &ld, piv, &info );
for( int j = 0; j < maxn; ++j ) {
for( int i = 0; i < j; ++i ) {
*LU(i,j) = *LU(j,i);
}
}
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
magma_ssetmatrix( m, m, LU, ld, dA, ld );
magma_ssetvector( m, C, 1, dC1, 1 );
magma_ssetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_strsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasStrsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasSaxpy( m, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_slange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_STRSM( MagmaLeft, m, 1 ) / 1e9;
printf( "strsv( %c, %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], diag[id], error, gflops/t1, gflops/t2 );
}}}
printf( "\n" );
printf( "========== Level 3 BLAS ==========\n" );
// ----- test SGEMM
// C = alpha*A*B + beta*C, with A m*k or k*m; B k*n or n*k; C m*n
// try combinations of no-trans/trans
for( int ia = 0; ia < 3; ++ia ) {
for( int ib = 0; ib < 3; ++ib ) {
bool nta = (trans[ia] == 'N');
bool ntb = (trans[ib] == 'N');
magma_ssetmatrix( (nta ? m : k), (nta ? m : k), A, ld, dA, ld );
magma_ssetmatrix( (ntb ? k : n), (ntb ? n : k), B, ld, dB, ld );
magma_ssetmatrix( m, n, C, ld, dC1, ld );
magma_ssetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_sgemm( trans[ia], trans[ib], m, n, k, alpha, dA, ld, dB, ld, beta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasSgemm( trans[ia], trans[ib], m, n, k, alpha, dA, ld, dB, ld, beta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasSaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_slange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_SGEMM( m, n, k ) / 1e9;
printf( "sgemm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
trans[ia], trans[ib], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test SSYMM
// C = alpha*A*B + beta*C (left) with A m*m symmetric; B,C m*n; or
// C = alpha*B*A + beta*C (right) with A n*n symmetric; B,C m*n
// try left/right, upper/lower
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
magma_ssetmatrix( m, m, A, ld, dA, ld );
magma_ssetmatrix( m, n, B, ld, dB, ld );
magma_ssetmatrix( m, n, C, ld, dC1, ld );
magma_ssetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ssymm( side[is], uplo[iu], m, n, alpha, dA, ld, dB, ld, beta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasSsymm( side[is], uplo[iu], m, n, alpha, dA, ld, dB, ld, beta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasSaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_slange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_SSYMM( side[is], m, n ) / 1e9;
printf( "ssymm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
side[is], uplo[iu], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test SSYRK
// C = alpha*A*A^H + beta*C (no-trans) with A m*k and C m*m symmetric; or
// C = alpha*A^H*A + beta*C (trans) with A k*m and C m*m symmetric
// try upper/lower, no-trans/trans
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
magma_ssetmatrix( n, k, A, ld, dA, ld );
magma_ssetmatrix( n, n, C, ld, dC1, ld );
magma_ssetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ssyrk( uplo[iu], trans[it], n, k, dalpha, dA, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasSsyrk( uplo[iu], trans[it], n, k, dalpha, dA, ld, dbeta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasSaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_slange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_SSYRK( k, n ) / 1e9;
printf( "ssyrk( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test SSYR2K
// C = alpha*A*B^H + ^alpha*B*A^H + beta*C (no-trans) with A,B n*k; C n*n symmetric; or
// C = alpha*A^H*B + ^alpha*B^H*A + beta*C (trans) with A,B k*n; C n*n symmetric
// try upper/lower, no-trans/trans
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
bool nt = (trans[it] == 'N');
magma_ssetmatrix( (nt ? n : k), (nt ? n : k), A, ld, dA, ld );
magma_ssetmatrix( n, n, C, ld, dC1, ld );
magma_ssetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ssyr2k( uplo[iu], trans[it], n, k, alpha, dA, ld, dB, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasSsyr2k( uplo[iu], trans[it], n, k, alpha, dA, ld, dB, ld, dbeta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasSaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_slange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_SSYR2K( k, n ) / 1e9;
printf( "ssyr2k( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test STRMM
// C = alpha*A*C (left) with A m*m triangular; C m*n; or
// C = alpha*C*A (right) with A n*n triangular; C m*n
// try left/right, upper/lower, no-trans/trans, unit/non-unit
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
bool left = (side[is] == 'L');
magma_ssetmatrix( (left ? m : n), (left ? m : n), A, ld, dA, ld );
magma_ssetmatrix( m, n, C, ld, dC1, ld );
magma_ssetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_strmm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasStrmm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasSaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_slange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_STRMM( side[is], m, n ) / 1e9;
printf( "strmm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// ----- test STRSM
// solve A*X = alpha*B (left) with A m*m triangular; B m*n; or
// solve X*A = alpha*B (right) with A n*n triangular; B m*n
// try left/right, upper/lower, no-trans/trans, unit/non-unit
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
bool left = (side[is] == 'L');
magma_ssetmatrix( (left ? m : n), (left ? m : n), LU, ld, dA, ld );
magma_ssetmatrix( m, n, C, ld, dC1, ld );
magma_ssetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_strsm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasStrsm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasSaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_sgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_slange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_STRSM( side[is], m, n ) / 1e9;
printf( "strsm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// cleanup
magma_free_cpu( piv );
magma_free_pinned( A );
magma_free_pinned( B );
magma_free_pinned( C );
magma_free_pinned( C2 );
magma_free_pinned( LU );
magma_free( dA );
magma_free( dB );
magma_free( dC1 );
magma_free( dC2 );
}
if ( total_error != 0. ) {
printf( "total error %.2g -- ought to be 0 -- some test failed (see above).\n",
total_error );
}
else {
printf( "all tests passed\n" );
}
TESTING_FINALIZE();
return 0;
}