gemm_impl.cpp 21.9 KB
Newer Older
1
2
3
/*
 * The MIT License (MIT)
 *
4
 * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved.
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
 *
 * Permission is hereby granted, free of charge, to any person obtaining a copy
 * of this software and associated documentation files (the "Software"), to deal
 * in the Software without restriction, including without limitation the rights
 * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
 * copies of the Software, and to permit persons to whom the Software is
 * furnished to do so, subject to the following conditions:
 *
 * The above copyright notice and this permission notice shall be included in
 * all copies or substantial portions of the Software.
 *
 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
 * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 * THE SOFTWARE.
 */
24

25
#include <rocblas/rocblas.h>
26
#include <migraphx/gpu/gemm_impl.hpp>
Paul's avatar
Paul committed
27
28
#include <migraphx/reduce_dims.hpp>
#include <migraphx/generate.hpp>
29
30
31
32
#include <migraphx/time.hpp>

using microseconds = std::chrono::duration<double, std::micro>;

Shucai Xiao's avatar
Shucai Xiao committed
33
34
35
36
namespace migraphx {
inline namespace MIGRAPHX_INLINE_NS {
namespace gpu {

37
// Convert rocBLAS datatypes to equivalent Migraphx data types
38
rocblas_datatype get_type(shape::type_t type)
Shucai Xiao's avatar
Shucai Xiao committed
39
{
40
    switch(type)
41
    {
42
43
44
45
46
47
48
    case shape::double_type: return rocblas_datatype_f64_r;
    case shape::float_type: return rocblas_datatype_f32_r;
    case shape::half_type: return rocblas_datatype_f16_r;
    case shape::int8_type: return rocblas_datatype_i8_r;
    case shape::uint8_type: return rocblas_datatype_u8_r;
    case shape::int32_type: return rocblas_datatype_i32_r;
    case shape::uint32_type: return rocblas_datatype_u32_r;
Paul Fultz II's avatar
Paul Fultz II committed
49
    case shape::tuple_type:
50
    case shape::bool_type:
51
52
53
54
    case shape::uint16_type:
    case shape::int16_type:
    case shape::int64_type:
    case shape::uint64_type: MIGRAPHX_THROW("ROCBLAS_GEMM: data type not supported!");
55
    }
56
57

    MIGRAPHX_THROW("ROCBLAS_GEMM: data type not supported!");
58
59
}

60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
void blas_shape(const shape& s)
{
    if(s.lens().size() < 2)
        return;
    if(std::none_of(s.strides().end() - 2, s.strides().end(), [&](auto i) { return i == 1; }))
        MIGRAPHX_THROW("GPU_GEMM: needs to have one matrix stride as 1");
    if(s.lens().size() < 3)
        return;
    shape batch_shape{s.type(),
                      {s.lens().begin(), s.lens().end() - 2},
                      {s.strides().begin(), s.strides().end() - 2}};
    auto batch_shapes = reduce_dims({batch_shape});
    if(batch_shapes.front().lens().size() != 1)
        MIGRAPHX_THROW("GPU_GEMM: Batch dimension is not collapsible");
}

76
77
78
79
80
81
82
83
84
85
86
87
88
shape transpose_batch(const shape& s, unsigned trans_batch)
{
    if(trans_batch == 0)
        return s;
    if(s.lens().size() < 3)
        return s;
    auto batch = s.lens().size() - 3;
    std::vector<int64_t> perm(s.lens().size());
    std::iota(perm.begin(), perm.end(), 0);
    std::swap(perm[batch], perm[batch + trans_batch]);
    return shape::from_permutation(s.type(), s.lens(), perm);
}

89
90
91
92
93
94
95
96
/**
 * Returns results of rocblas_status_success, rocblas_status_perf_degraded,
 * or rocblas_status_invalid_value.  Caller
 * is expected to check for invalid index.  Any other result causes an exception.
 *
 */
template <class F, class Pack, class... Ts>
auto rocblas_invoke(F f, Pack p, Ts... xs)
97
{
98
99
100
101
102
103
104
105
106
107
108
109
110
111
    return p([=](auto... ws) {
        auto status = f(ws..., xs...);
        if(status != rocblas_status_success and status != rocblas_status_invalid_value)
        {
            if(status == rocblas_status_perf_degraded)
            {
                std::cerr << "WARNING: degraded perf. in rocBLAS call" << std::endl;
            }
            else
                MIGRAPHX_THROW("rocblas_invoke: rocBLAS call failed with status " +
                               std::to_string(status));
        }
        return status;
    });
112
113
}

114
static bool is_transposed(const shape& s) { return s.transposed() and s.strides().back() != 1; }
115

116
static rocblas_int get_batch_stride(const shape& s)
117
{
118
    // This value is not needed for non-strided inputs
119
    if(s.strides().size() < 3)
120
121
        return 0;
    else
122
        return s.strides()[s.strides().size() - 3];
123
124
}

125
126
127
128
129
/**
 * Wrapper for multiple rocBLAS calls.  The constructor creates parameters for
 * these calls based on data shapes and other values contained in the associated
 * instruction and operation.
 *
130
 * The template parameter T is not the type of the matrix data but of the weighting
131
132
133
134
 * coefficients alpha and beta (these are float in rocBLAS internals)
 */
template <typename T>
struct gemm_impl
Shucai Xiao's avatar
Shucai Xiao committed
135
{
136
137
138
139
140
141
142
143
144
    gemm_impl(const shape& output_shape,
              const std::vector<shape>& input_shapes,
              T alpha_param,
              T beta_param,
              bool compute_fp32_flag)
        : alpha(alpha_param),
          beta(beta_param),
          is_3inputs(input_shapes.size() == 4),
          compute_fp32(compute_fp32_flag)
145
    {
146
147
148
149
        if(not is_3inputs)
        {
            beta = 0;
        }
Paul's avatar
Format  
Paul committed
150
151
152
153
154
155
156

        // Create lambdas that will cast alpha, beta to the output shape's type
        // and retain the values being pointed to
        output_shape.visit_type([&](auto as) {
            auto alpha_r = as(alpha);
            auto beta_r  = as(beta);
            if(compute_fp32)
157
            {
Paul's avatar
Format  
Paul committed
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
                get_alpha = [=] { return &alpha; };
                get_beta  = [=] { return &beta; };
            }
            else
            {
                get_alpha = [=] { return &alpha_r; };
                get_beta  = [=] { return &beta_r; };
            }
        });

        transa     = is_transposed(input_shapes[0]);
        transb     = is_transposed(input_shapes[1]);
        auto n_dim = output_shape.lens().size();
        auto dim_0 = n_dim - 2;
        auto dim_1 = n_dim - 1;
        // Leading dimensions of matrices
        lda = input_shapes[0].strides()[transa ? dim_1 : dim_0];
        ldb = input_shapes[1].strides()[transb ? dim_1 : dim_0];
        ldc = input_shapes[2].strides()[dim_0];
        ldd = is_3inputs ? input_shapes[3].strides()[dim_0] : ldc;

        arg_type    = get_type(input_shapes[0].type());
        output_type = arg_type;
        if(output_type == rocblas_datatype_i8_r)
        {
            output_type = rocblas_datatype_i32_r;
Paul's avatar
Format  
Paul committed
184
        }
Paul's avatar
Paul committed
185
186
187
        compute_type = output_type;
        if(compute_fp32)
        {
Paul's avatar
Format  
Paul committed
188
189
            if(arg_type == rocblas_datatype_f16_r)
                compute_type = rocblas_datatype_f32_r;
Paul's avatar
Paul committed
190
        }
191

Paul's avatar
Paul committed
192
193
        auto a_lens = input_shapes[0].lens();
        auto b_lens = input_shapes[1].lens();
194

Paul's avatar
Paul committed
195
196
197
198
        auto out_lens = output_shape.lens();
        m             = out_lens[dim_0];
        n             = out_lens[dim_1];
        k             = input_shapes[0].lens()[dim_1];
Shucai Xiao's avatar
Shucai Xiao committed
199

Paul's avatar
Paul committed
200
201
202
203
        a_stride     = get_batch_stride(input_shapes[0]);
        b_stride     = get_batch_stride(input_shapes[1]);
        c_stride     = get_batch_stride(input_shapes[2]);
        d_stride     = is_3inputs ? get_batch_stride(input_shapes[3]) : c_stride;
Paul's avatar
Format  
Paul committed
204
205
        num_matrices = std::accumulate(
            out_lens.rbegin() + 2, out_lens.rend(), std::size_t{1}, std::multiplies<std::size_t>());
Paul's avatar
Paul committed
206
207
        if(num_matrices == 1 or (num_matrices > 1 and b_stride == 0))
        {
Paul's avatar
Format  
Paul committed
208
209
210
211
212
            // If the batch dimension of B is broadcasted, then we can
            // multiply m by the batch_size and use rocblas_gemm_ex
            // instead of rocblas_gemm_strided_batched_ex.
            m *= num_matrices;
            strided_batched = false;
Paul's avatar
Paul committed
213
        }
214
    }
215

216
217
218
219
220
221
222
    void run(context& ctx, const std::vector<argument>& input_args, int32_t solution_idx = 0) const
    {
        if(strided_batched)
        {
            auto common_args = create_strided_batched_args_common(ctx, input_args);
            rocblas_invoke(&rocblas_gemm_strided_batched_ex,
                           common_args,
Paul's avatar
Paul committed
223
                           rocblas_gemm_algo_solution_index,
224
                           solution_idx,
Paul's avatar
Paul committed
225
                           gemm_flags);
Shucai Xiao's avatar
Shucai Xiao committed
226
227
228
        }
        else
        {
229
            auto common_args = create_gemm_ex_args_common(ctx, input_args);
Paul's avatar
Format  
Paul committed
230
231
            rocblas_invoke(&rocblas_gemm_ex,
                           common_args,
Paul's avatar
Paul committed
232
                           rocblas_gemm_algo_solution_index,
Paul's avatar
Format  
Paul committed
233
234
                           solution_idx,
                           gemm_flags);
Shucai Xiao's avatar
Shucai Xiao committed
235
        }
236
237
    }

Paul's avatar
Paul committed
238
#ifdef MIGRAPHX_USE_ROCBLAS_TUNING_API
239
240
    auto validate(context& ctx, const std::vector<shape>& input_shapes, int32_t solution_idx) const
    {
241
        // Create dummy arguments for the shapes, and call the overloaded method
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
        std::vector<argument> input_args;
        std::transform(input_shapes.begin(),
                       input_shapes.end(),
                       std::back_inserter(input_args),
                       [](const shape& x) { return to_gpu(generate_argument(x)); });

        return validate(ctx, input_args, solution_idx);
    }

    /**
     * Checks a particular solution for validity by running it with the flag
     * rocblas_gemm_flags_check_solution_index (could be invalid if this model was
     * tuned with a different rocBLAS version)
     *
     * @return Returns either solution_idx if valid, or else the default value 0
     * if not.  The default does not mean list index 0, but tells the picker
     * to choose a solution.
     */
    int32_t
    validate(context& ctx, const std::vector<argument>& input_args, int32_t solution_idx) const
    {
        rocblas_status_ check_valid(rocblas_status_success);

        if(strided_batched)
        {
            auto common_args = create_strided_batched_args_common(ctx, input_args);
            check_valid      = rocblas_invoke(&rocblas_gemm_strided_batched_ex,
                                         common_args,
                                         rocblas_gemm_algo_solution_index,
                                         solution_idx,
                                         rocblas_gemm_flags_check_solution_index);
        }
        else
        {
            auto common_args = create_gemm_ex_args_common(ctx, input_args);
            check_valid      = rocblas_invoke(&rocblas_gemm_ex,
                                         common_args,
279
                                         rocblas_gemm_algo_solution_index,
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
                                         solution_idx,
                                         rocblas_gemm_flags_check_solution_index);
        }

        if(check_valid == rocblas_status_invalid_value)
        {
            std::cerr << "WARNING:  tuned solution is invalid; reverting to default" << std::endl;
            return 0;
        }
        return solution_idx;
    }
#endif

    /**
     * Helper method to create that subset of a long rocBLAS argument list that is common
     * to multiple "...strided_batched..." calls.
     *
     * The rocblas_gemm API handles inputs and output matrices as
     *  column-major format. When doing a C = A * B, we actually do
     *  C^T = (B^T) * (A^T). That is the reason we input args[1] as
     *   A and args[0] as B in calling the rocblas_gemm.
     *
     */
    auto create_strided_batched_args_common(context& ctx, const std::vector<argument>& args) const
    {
        return pack(ctx.get_stream().get_rocblas(),
                    transb ? rocblas_operation_transpose : rocblas_operation_none,
                    transa ? rocblas_operation_transpose : rocblas_operation_none,
                    n,
                    m,
                    k,
                    get_alpha(),
                    args[1].data(),
                    arg_type,
                    ldb,
                    b_stride,
                    args[0].data(),
                    arg_type,
                    lda,
                    a_stride,
                    get_beta(),
                    args[2].data(),
                    output_type,
                    ldc,
                    c_stride,
                    is_3inputs ? args[3].data() : args[2].data(),
                    output_type,
                    ldd,
                    d_stride,
                    num_matrices,
                    compute_type);
    }

    /**
     * Helper method to create that subset of a long rocBLAS argument list that is common
     * to multiple "gemm_ex..." calls.
     *
     * The rocblas_gemm API handles inputs and output matrices as
     *  column-major format. When doing a C = A * B, we actually do
     *   C^T = (B^T) * (A^T). That is the reason we input args[1] as
     *   A and args[0] as B in calling the rocblas_gemm.
     *
     * */
    auto create_gemm_ex_args_common(context& ctx, const std::vector<argument>& args) const
    {
        return pack(ctx.get_stream().get_rocblas(),
                    transb ? rocblas_operation_transpose : rocblas_operation_none,
                    transa ? rocblas_operation_transpose : rocblas_operation_none,
                    n,
                    m,
                    k,
                    get_alpha(),
                    args[1].data(),
                    arg_type,
                    ldb,
                    args[0].data(),
                    arg_type,
                    lda,
                    get_beta(),
                    args[2].data(),
                    output_type,
                    ldc,
                    is_3inputs ? args[3].data() : args[2].data(),
                    output_type,
                    ldd,
                    compute_type);
    }
Paul's avatar
Paul committed
367
#ifdef MIGRAPHX_USE_ROCBLAS_TUNING_API
368
369
370
371
372
373
374
    /**
     * Find best rocBLAS solution:  Get list of solutions and try them all, returning the index
     * of the fastest one.
     */
    int tune(context& ctx, const std::vector<shape>& input_shapes) const
    {
        // tuning meta parameters
375
        const int hot_calls = 40;
376
377
378
379
380
381
382
383
384
385
386
387
388
389
390
391
392
393
394

        std::vector<argument> input_args;
        std::transform(input_shapes.begin(),
                       input_shapes.end(),
                       std::back_inserter(input_args),
                       [](const shape& x) { return to_gpu(generate_argument(x)); });

        // Get the solutions list in 2 rocBLAS steps:
        // 1.  Find out how many solutions there are and allocate the array
        // 2.  Get the solutions
        //
        rocblas_int list_size = 0;
        std::vector<rocblas_int> solution_indices;
        if(strided_batched)
        {
            auto common_args = create_strided_batched_args_common(ctx, input_args);
            rocblas_invoke(&rocblas_gemm_strided_batched_ex_get_solutions,
                           common_args,
                           rocblas_gemm_algo_solution_index,
Paul's avatar
Paul committed
395
                           gemm_flags,
396
397
398
399
400
401
402
403
                           nullptr,
                           &list_size);
            solution_indices.resize(list_size);

            auto common_sol_args = create_strided_batched_args_common(ctx, input_args);
            rocblas_invoke(&rocblas_gemm_strided_batched_ex_get_solutions,
                           common_sol_args,
                           rocblas_gemm_algo_solution_index,
Paul's avatar
Paul committed
404
                           gemm_flags,
405
406
407
408
409
410
411
412
413
                           solution_indices.data(),
                           &list_size);
        }
        else
        {
            auto common_args = create_gemm_ex_args_common(ctx, input_args);
            rocblas_invoke(&rocblas_gemm_ex_get_solutions,
                           common_args,
                           rocblas_gemm_algo_solution_index,
Paul's avatar
Paul committed
414
                           gemm_flags,
415
416
417
418
419
420
421
422
                           nullptr,
                           &list_size);
            solution_indices.resize(list_size);

            auto common_sol_args = create_gemm_ex_args_common(ctx, input_args);
            rocblas_invoke(&rocblas_gemm_ex_get_solutions,
                           common_sol_args,
                           rocblas_gemm_algo_solution_index,
Paul's avatar
Paul committed
423
                           gemm_flags,
424
425
426
427
                           solution_indices.data(),
                           &list_size);
        }

428
        double best_time  = std::numeric_limits<double>::max();
429
430
        double first_time = -1;
        // Initialize to default solution index
431
        rocblas_int best_sol = 0;
432
433
        for(auto sol : solution_indices)
        {
434
435
            // Warmup: the first call to an op. may not be representative since there is
            // more time taken initializing caches, etc. so we won't time it.
Paul's avatar
Paul committed
436
437
            run(ctx, input_args, sol);
            double host_time = time<milliseconds>([&] {
Paul's avatar
Format  
Paul committed
438
                for([[maybe_unused]] int hc : range(hot_calls))
Paul's avatar
Paul committed
439
                    run(ctx, input_args, sol);
440
                ctx.finish();
Paul's avatar
Paul committed
441
442
            });

443
444
            host_time /= hot_calls;

445
            // dev/evaluation only: track time for first solution.
446
447
448
449
            if(first_time < 0)
                first_time = host_time;

            // track current best
450
            if(host_time < best_time)
451
            {
452
453
                best_sol  = sol;
                best_time = host_time;
454
455
            }
        }
Paul's avatar
Paul committed
456
457
        std::cout << "Winning GEMM solution: " << best_sol << " in " << best_time << " ms, beats "
                  << first_time << "ms" << std::endl;
458
        return best_sol;
459
460
461
    }
#endif
    private:
Paul's avatar
Paul committed
462
    size_t num_matrices = 0;
Paul's avatar
Format  
Paul committed
463
464
465
466
467
468
469
    rocblas_int m       = 0;
    rocblas_int n       = 0;
    rocblas_int k       = 0;
    bool transa         = false;
    bool transb         = false;
    T alpha             = 0;
    T beta              = 0;
Paul's avatar
Paul committed
470
471
472

    std::function<const void*()> get_alpha{};
    std::function<const void*()> get_beta{};
Paul's avatar
Format  
Paul committed
473
    rocblas_gemm_flags gemm_flags = rocblas_gemm_flags_none;
Paul's avatar
Format  
Paul committed
474
475
476
477
478
479
480
481
    rocblas_int lda               = 0;
    rocblas_int ldb               = 0;
    rocblas_int ldc               = 0;
    rocblas_int ldd               = 0;
    rocblas_int a_stride          = 0;
    rocblas_int b_stride          = 0;
    rocblas_int c_stride          = 0;
    rocblas_int d_stride          = 0;
Paul's avatar
Paul committed
482
    rocblas_datatype compute_type = rocblas_datatype_f32_r;
Paul's avatar
Format  
Paul committed
483
484
485
486
487
    rocblas_datatype arg_type     = rocblas_datatype_f32_r;
    rocblas_datatype output_type  = rocblas_datatype_f32_r;
    bool strided_batched          = true;
    bool is_3inputs               = true;
    bool compute_fp32             = true;
488
489
490
491
492
493
494
495
496
497
498
499
500
501
502
}; // gemm_impl

void gemm_compute(context& ctx,
                  const shape& output_shape,
                  const std::vector<argument>& args,
                  float alpha,
                  float beta,
                  bool compute_fp32,
                  int32_t solution_idx)
{
    std::vector<shape> input_shapes;
    std::transform(args.begin(),
                   args.end(),
                   std::back_inserter(input_shapes),
                   [](const argument& x) { return x.get_shape(); });
Paul's avatar
Format  
Paul committed
503
    auto gemm_item = gemm_impl<float>(output_shape, input_shapes, alpha, beta, compute_fp32);
504
505
506
507
508
509
510
511
512
513
514
515
516
517
518
519
    gemm_item.run(ctx, args, solution_idx);
}

void gemm_compute(context& ctx,
                  const shape& output_shape,
                  const std::vector<argument>& args,
                  int32_t alpha,
                  int32_t beta,
                  bool compute_fp32,
                  int32_t solution_idx)
{
    std::vector<shape> input_shapes;
    std::transform(args.begin(),
                   args.end(),
                   std::back_inserter(input_shapes),
                   [](const argument& x) { return x.get_shape(); });
Paul's avatar
Format  
Paul committed
520
    auto gemm_item = gemm_impl<int32_t>(output_shape, input_shapes, alpha, beta, compute_fp32);
521
    gemm_item.run(ctx, args, solution_idx);
522
}
Shucai Xiao's avatar
Shucai Xiao committed
523

524
525
526
527
/**
 * Decides if the tune() or validate() method is appropriate and calls it.
 * Return value is the chosen solution index, or 0 to let picker choose it.
 */
528
529
530
531
532
533
534
int32_t gemm_finalize(context& ctx,
                      const shape& output_shape,
                      const std::vector<shape>& input_shapes,
                      float alpha,
                      float beta,
                      bool compute_fp32,
                      int32_t solution_idx)
535
{
Paul's avatar
Paul committed
536
#ifdef MIGRAPHX_USE_ROCBLAS_TUNING_API
537
538
539
540
541

    // This code should be called only if either the environment var.
    // MIGRAPHX_ENABLE_GEMM_TUNING, or option --exhaustive-tune, is set

    if(solution_idx == 0)
542
    {
Paul's avatar
Format  
Paul committed
543
        auto gemm_item = gemm_impl<float>(output_shape, input_shapes, alpha, beta, compute_fp32);
544
545
        solution_idx = gemm_item.tune(ctx, input_shapes);
    }
546
    else
547
548
549
    {
        // If a tuned solution index is already given, don't tune again but validate
        // in case the data was tuned with a different rocBLAS version
Paul's avatar
Format  
Paul committed
550
        auto gemm_item = gemm_impl<float>(output_shape, input_shapes, alpha, beta, compute_fp32);
551
552
553
554
        solution_idx = gemm_item.validate(ctx, input_shapes, solution_idx);
    }
#else
    (void)ctx, (void)output_shape, (void)input_shapes;
Paul's avatar
Paul committed
555
    (void)alpha, (void)beta, (void)compute_fp32;
556
557
#endif
    return solution_idx;
558
559
}

560
561
/**
 * Decides if the tune() or validate() method is appropriate and calls it.
562
 * Return value is the chosen solution index, or 0 to let picker choose it.
563
 */
564
565
566
567
568
569
570
int32_t gemm_finalize(context& ctx,
                      const shape& output_shape,
                      const std::vector<shape>& input_shapes,
                      int32_t alpha,
                      int32_t beta,
                      bool compute_fp32,
                      int32_t solution_idx)
571
{
Paul's avatar
Paul committed
572
#ifdef MIGRAPHX_USE_ROCBLAS_TUNING_API
573
    if(solution_idx == 0)
574
    {
Paul's avatar
Format  
Paul committed
575
        auto gemm_item = gemm_impl<int32_t>(output_shape, input_shapes, alpha, beta, compute_fp32);
576
577
        solution_idx = gemm_item.tune(ctx, input_shapes);
    }
578
    else
579
580
581
    {
        // If a tuned solution index is already given, don't tune again but validate
        // in case the data was tuned with a different rocBLAS version
Paul's avatar
Format  
Paul committed
582
        auto gemm_item = gemm_impl<int32_t>(output_shape, input_shapes, alpha, beta, compute_fp32);
583
584
585
586
        solution_idx = gemm_item.validate(ctx, input_shapes, solution_idx);
    }
#else
    (void)ctx, (void)output_shape, (void)input_shapes;
Paul's avatar
Paul committed
587
    (void)alpha, (void)beta, (void)compute_fp32;
588
589
#endif
    return solution_idx;
Shucai Xiao's avatar
Shucai Xiao committed
590
591
592
593
594
}

} // namespace gpu
} // namespace MIGRAPHX_INLINE_NS
} // namespace migraphx