* Tests the automatic GEMM tuning feature for rocBLAS and hipBLASLt. * In the finalize() method of the gemm op, * rocBLAS API functions are called to quickly benchmark all the GEMM solutions * available in the currently installed rocBLAS library and choose the index of the fastest. */
| 54 | * available in the currently installed rocBLAS library and choose the index of the fastest. |
| 55 | */ |
| 56 | TEST_CASE(gemm_tune) |
| 57 | { |
| 58 | migraphx::program p; |
| 59 | auto* mm = p.get_main_module(); |
| 60 | |
| 61 | migraphx::shape sa{migraphx::shape::float_type, {4, 2}}; |
| 62 | migraphx::shape sb{migraphx::shape::float_type, {2, 3}}; |
| 63 | auto a = mm->add_parameter("a", sa); |
| 64 | auto b = mm->add_parameter("b", sb); |
| 65 | |
| 66 | migraphx::operation dot_op = migraphx::make_op("dot"); |
| 67 | mm->add_instruction(dot_op, a, b); |
| 68 | |
| 69 | // lowering adds gemm implementation for dot operator |
| 70 | run_lowering(p); |
| 71 | |
| 72 | migraphx::target gpu_t = migraphx::gpu::target{}; |
| 73 | migraphx::compile_options options; |
| 74 | options.exhaustive_tune = true; |
| 75 | p.compile(gpu_t, options); |
| 76 | |
| 77 | migraphx::value solution_idx(0); |
| 78 | for(auto ins : iterator_for(*p.get_main_module())) |
| 79 | { |
| 80 | if(ins->name() == "gpu::gemm" or ins->name() == "gpu::hip_gemm") |
| 81 | { |
| 82 | auto gemm_op = migraphx::get_operation(ins); |
| 83 | |
| 84 | // tuned solution index is not deterministic, but anything other than 0 |
| 85 | // (default, invalid, or not available) is good. |
| 86 | // gemm_op.to_value().debug_print(); |
| 87 | solution_idx = gemm_op.to_value()["solution_idx"]; |
| 88 | break; |
| 89 | } |
| 90 | } |
| 91 | #if defined(MIGRAPHX_USE_ROCBLAS_TUNING_API) or MIGRAPHX_USE_HIPBLASLT |
| 92 | EXPECT(0 != solution_idx.to<std::size_t>()); |
| 93 | #else |
| 94 | EXPECT(0 == solution_idx.to<std::size_t>()); |
| 95 | #endif |
| 96 | } |
| 97 | |
| 98 | // GEMM tuning of a strided-batch matrix; invokes rocblas_gemm_strided_batched_ex |
| 99 | TEST_CASE(gemm_tune_strided) |
nothing calls this directly
no test coverage detected