diff --git a/projects/miopen/docs/how-to/debug-log.rst b/projects/miopen/docs/how-to/debug-log.rst index 10297c806c42..0d7cc67dbc2f 100644 --- a/projects/miopen/docs/how-to/debug-log.rst +++ b/projects/miopen/docs/how-to/debug-log.rst @@ -151,7 +151,7 @@ Some of the solutions have individual controls, which affect both find and immed * ``MIOPEN_DEBUG_CONV_DIRECT_HIP_FWD11X11`` -- ``ConvHipDirectFwd11x11`` * ``MIOPEN_DEBUG_CONV_DIRECT_OCL_FWDGEN`` -- ``ConvOclDirectFwdGen`` * ``MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD`` -- ``ConvOclDirectFwd`` - * ``MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1`` -- ``ConvOclDirectFwd1x1`` + * ``MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW2`` -- ``ConvOclBwdWrW2`` (where n = ``{1,2,4,8,16}``) and ``ConvOclBwdWrW2NonTunable`` * ``MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW53`` -- ``ConvOclBwdWrW53`` diff --git a/projects/miopen/docs/reference/env_variables.rst b/projects/miopen/docs/reference/env_variables.rst index 9f279b9574c1..8f3c804534de 100644 --- a/projects/miopen/docs/reference/env_variables.rst +++ b/projects/miopen/docs/reference/env_variables.rst @@ -278,11 +278,6 @@ For more information, see :doc:`Logging and debugging <../how-to/debug-log>`. - | 0: Disable | 1: Enable - * - | ``MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1`` - | Controls ConvOclDirectFwd1x1 direct solution. - - | 0: Disable - | 1: Enable - * - | ``MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW2`` | Controls ConvOclBwdWrW2 (n={1,2,4,8,16}) and ConvOclBwdWrW2NonTunable solutions. - | 0: Disable diff --git a/projects/miopen/src/CMakeLists.txt b/projects/miopen/src/CMakeLists.txt index 00429d7b4bd6..9e6f401b7506 100644 --- a/projects/miopen/src/CMakeLists.txt +++ b/projects/miopen/src/CMakeLists.txt @@ -204,6 +204,7 @@ set( MIOpen_Source solver/conv/conv_direct_naive_conv_bwd.cpp solver/conv/conv_direct_naive_conv_fwd.cpp solver/conv/conv_direct_naive_conv_wrw.cpp + solver/conv/conv_hip_implicit_gemm_bwd_v1r1.cpp solver/conv/conv_hip_implicit_gemm_bwd_v1r1_xdlops.cpp solver/conv/conv_hip_implicit_gemm_bwd_v4r1.cpp @@ -238,7 +239,6 @@ set( MIOpen_Source solver/conv/conv_hip_dir2D11x11.cpp solver/conv/conv_ocl_dir2Dfwd.cpp solver/conv/conv_ocl_dir2Dfwd_exhaustive_search.cpp - solver/conv/conv_ocl_dir2Dfwd1x1.cpp solver/conv/conv_ocl_dir2Dfwdgen.cpp solver/conv/conv_wino_fury_RxS.cpp solver/conv/conv_wino_rage_RxS.cpp @@ -566,9 +566,7 @@ if( MIOPEN_BACKEND MATCHES "OpenCL" OR MIOPEN_BACKEND STREQUAL "HIPOC" OR MIOPEN kernels/MIOpenPoolingBwdND.cpp kernels/MIOpenPoolingForwardNaive.cpp kernels/MIOpenPoolingND.cpp - kernels/MIOpenConv1x1S.cl - kernels/MIOpenConv1x1J1.cl - kernels/MIOpenConv1x1J1_stride.cl + kernels/MIOpenReduceCalculation.cpp kernels/MIOpenReduceExtreme.cpp kernels/MIOpenReduceSum.cpp diff --git a/projects/miopen/src/fin/fin_interface.cpp b/projects/miopen/src/fin/fin_interface.cpp index 5298bd57390a..9c74871a016e 100644 --- a/projects/miopen/src/fin/fin_interface.cpp +++ b/projects/miopen/src/fin/fin_interface.cpp @@ -1,28 +1,5 @@ -/******************************************************************************* - * - * MIT License - * - * Copyright (c) 2024 Advanced Micro Devices, Inc. - * - * 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. - * - *******************************************************************************/ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT #include @@ -286,7 +263,6 @@ AnySolver::AnySolver case 2: SetObject(); break; case 3: SetObject(); break; case 11: SetObject(); break; - case 13: SetObject(); break; case 16: SetObject(); break; case 17: SetObject(); break; case 18: SetObject>(); break; diff --git a/projects/miopen/src/include/miopen/conv/solvers.hpp b/projects/miopen/src/include/miopen/conv/solvers.hpp index 49f24fa374fc..e27ea412219f 100644 --- a/projects/miopen/src/include/miopen/conv/solvers.hpp +++ b/projects/miopen/src/include/miopen/conv/solvers.hpp @@ -1,28 +1,5 @@ -/******************************************************************************* - * - * MIT License - * - * Copyright (c) 2026 Advanced Micro Devices, Inc. - * - * 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. - * - *******************************************************************************/ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT #pragma once @@ -1719,24 +1696,6 @@ struct MIOPEN_INTERNALS_EXPORT ConvOclDirectFwd final : ConvOclDirectFwdLegacyEx const LegacyPerformanceConfig&) const override; }; -struct MIOPEN_INTERNALS_EXPORT ConvOclDirectFwd1x1 final : ConvOclDirectFwdLegacyExhaustiveSearch -{ - const std::string& SolverDbId() const override { return GetSolverDbId(); } - - bool IsApplicable(const ExecutionContext&, - const miopen::conv::ProblemDescription&) const override; - ConvSolution GetSolution(const ExecutionContext&, - const miopen::conv::ProblemDescription&, - const LegacyPerformanceConfig&) const override; - - bool IsValidPerformanceConfig(const ExecutionContext&, - const miopen::conv::ProblemDescription&, - const LegacyPerformanceConfig&) const override - { - return true; - } -}; - struct MIOPEN_INTERNALS_EXPORT ConvBinWinograd3x3U final : ConvSolver { const std::string& SolverDbId() const override { return GetSolverDbId(); } diff --git a/projects/miopen/src/kernels/MIOpenConv1x1J1.cl b/projects/miopen/src/kernels/MIOpenConv1x1J1.cl deleted file mode 100644 index 6b34d10bd10b..000000000000 --- a/projects/miopen/src/kernels/MIOpenConv1x1J1.cl +++ /dev/null @@ -1,401 +0,0 @@ -/******************************************************************************* - * - * MIT License - * - * Copyright (c) 2017 Advanced Micro Devices, Inc. - * - * 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. - * - *******************************************************************************/ - -// Trying to use float ATOMIC_ADD to increase total waves -// For example, 7x7 : 49 -// Global 0: H * W * N * (K/ 16 outplane per threads ) * (( Inputlanes/ 128 ) or 1) -// N * H * W = -// [groupId / ((K/ 16 outplane per threads ) * (( Inputlanes/ 128 ) or 1))] * 64 + localId -// For example, H7*W7*N*C256*K64 -// FIRST 4 waves -// Hit L2 for read -// try to Hit L2 for Atomic_ADD -// Assembly Shader can utlize LDS to Reductiion -// However Shader Compiler of OpenCL will compile any localId[1] to FLAT_BUFFER_LOAD not constant -// load - -// WAVE 0 == N0_7x7 + N1_(7x2+1), output 0-15 from K=64, Inputplanes from 0-127 -// WAVE 1 == N0_7x7 + N1_(7x2+1), output 16-31 from K=64, Inputplanes from 0-127 -// WAVE 2 == N0_7x7 + N1_(7x2+1), output 31-47 from K=64, Inputplanes from 0-127 -// WAVE 3 == N0_7x7 + N1_(7x2+1), output 48-63 from K=64, Inputplanes from 0-127 -// 2nd 4 waves -// Hit L2 for read -// try to Hit L2 for Atomic_ADD - -// WAVE 4 == N0_7x7 + N1_(7x2+1), output 0-15 from K=64, Inputplanes from 128-255 -// WAVE 5 == N0_7x7 + N1_(7x2+1), output 16-31 from K=64, Inputplanes from 128-255 -// WAVE 6 == N0_7x7 + N1_(7x2+1), output 31-47 from K=64, Inputplanes from 128-255 -// WAVE 7 == N0_7x7 + N1_(7x2+1), output 48-63 from K=64, Inputplanes from 128-255 - -// example -#if 0 // ndef MLopen_RUNNING -#define BATCHSIZE 32 -#define MLO_N_LCL_IN_MAPS_ONCE 8 - -#define H 28 -#define W 28 -#define C 192 -#define K 64 - -#define MLO_IN_HEIGHT H -#define MLO_IN_WIDTH W -#define MLO_N_INPUTS C - -//128 or MLO_N_INPUTS -#define MLO_N_LCL_IN_MAPS 192 - -#define MLO_N_OUTPUTS K - -#define H_out 28 -#define W_out 28 -#define MLO_N_LCL_OUT_MAPS 16 - -#define MLO_N_IN_GROUPS ((MLO_N_INPUTS + MLO_N_LCL_IN_MAPS - 1) / MLO_N_LCL_IN_MAPS) -#define MLO_CLOOP0 (MLO_N_LCL_IN_MAPS / MLO_N_LCL_IN_MAPS_ONCE) -#define MLO_CLOOP2 \ - ((MLO_N_INPUTS - MLO_N_LCL_IN_MAPS * (MLO_N_IN_GROUPS - 1)) / MLO_N_LCL_IN_MAPS_ONCE) -#define MLO_CHEAT_SHADER_COMPILER 1 - -#endif - -#define MLO_IN_CHANNEL_STRIDE (H * W) -#define MLO_IN_BATCH_STRIDE (H * W * C) - -#define MLO_WEI_BSTRIDE (1 * 1 * C * K) -#define MLO_WEI_CHANNEL_STRIDE (1 * 1 * C) - -#define MLO_OUT_BATCH_STRIDE (H_out * W_out * K) -#define MLO_OUT_CHANNEL_STRIDE (H_out * W_out) - -#define FIXED_WORKGROUP_SIZE 64 - -#define MLO_N_OUT_GROUPS (MLO_N_OUTPUTS / MLO_N_LCL_OUT_MAPS) - -#define MLO_GRP_SZ0 64 -#define MLO_GRP_SZ1 1 -#define MLO_GRP_SZ2 1 - -#define PPCAT_NX(A, B) A##B -#define PPCAT(A, B) PPCAT_NX(A, B) -#define TWO 2 -#define FOUR 4 -#define EIGHT 8 - -#if MIOPEN_USE_FP16 == 1 -#pragma OPENCL EXTENSION cl_khr_fp16 : enable -#define _FLOAT half -#define _UNION_FLOAT_T half2 -#define INIT(A) ((half2)(A[0], A[1])) -#ifndef HALF_MAX -#define MAX_VAL 65504 /* max value */ -#else -#define MAX_VAL HALF_MAX -#endif -#endif -#if MIOPEN_USE_FP32 == 1 -#define _FLOAT float -#define _UNION_FLOAT_T float -#define INIT(A) (A[0]) -#ifndef FLT_MAX -#define MAX_VAL 3.402823466e+38F /* max value */ -#else -#define MAX_VAL FLT_MAX -#endif -#endif - -#define _FLOAT2 PPCAT(_FLOAT, TWO) -#define _FLOAT4 PPCAT(_FLOAT, FOUR) -#define _FLOAT8 PPCAT(_FLOAT, EIGHT) - -#define MLO_CONV_BIAS 0 -#define UNUSED __attribute((__unused__)) - -typedef union -{ - unsigned int intVal; - _UNION_FLOAT_T floatVal; -} starVal; - -inline void AtomicAdd(volatile __global _FLOAT* source, const _FLOAT operand) -{ - starVal newVal, prevVal; - - prevVal.floatVal = INIT(source); - while(true) - { -#if MIOPEN_USE_FP16 == 1 - newVal.floatVal = (_FLOAT2)(prevVal.floatVal.x + operand, source[1]); -#endif -#if MIOPEN_USE_FP32 == 1 - newVal.floatVal = prevVal.floatVal + operand; -#endif - newVal.intVal = - atomic_cmpxchg((volatile __global unsigned int*)source, prevVal.intVal, newVal.intVal); - - // equal to pass - if(newVal.intVal == prevVal.intVal) - break; - - prevVal.intVal = newVal.intVal; - } -} - -__attribute__((reqd_work_group_size(MLO_GRP_SZ0, MLO_GRP_SZ1, MLO_GRP_SZ2))) __kernel void -MIOpenConv1x1(const __global _FLOAT* __restrict in_ptr, - __constant _FLOAT* __restrict wei_ptr, -#if MLO_CONV_BIAS - const __global _FLOAT* __restrict bias, -#endif - __global _FLOAT* __restrict out_ptr, - UNUSED _FLOAT dummy_val // nothing -) -{ - - uint grp_id0 = get_group_id(0); - uint out_grp_block = grp_id0 % MLO_N_OUT_GROUPS; - uint in_grp_block = (uint)(grp_id0 / MLO_N_OUT_GROUPS) % MLO_N_IN_GROUPS; - uint grp_id0_faked = (uint)(grp_id0 / MLO_N_OUT_GROUPS) / MLO_N_IN_GROUPS; - - uint local_id0 = get_local_id(0); -#if MLO_CHEAT_SHADER_COMPILER == 1 - uint grp_id2 = get_group_id(2); -#endif - - uint pos = (grp_id0_faked * FIXED_WORKGROUP_SIZE + local_id0) % MLO_IN_CHANNEL_STRIDE; - uint batch_id = (grp_id0_faked * FIXED_WORKGROUP_SIZE + local_id0) / MLO_IN_CHANNEL_STRIDE; - - if(batch_id >= BATCHSIZE) - return; - - uint out_id = out_grp_block * MLO_N_LCL_OUT_MAPS; - - uint gbl_in_off = batch_id * MLO_IN_BATCH_STRIDE + - in_grp_block * MLO_N_LCL_IN_MAPS * MLO_IN_CHANNEL_STRIDE + pos; - - uint wei_off = out_id * MLO_WEI_CHANNEL_STRIDE + in_grp_block * MLO_N_LCL_IN_MAPS; - - _FLOAT accum[MLO_N_LCL_OUT_MAPS]; - _FLOAT weights[MLO_N_LCL_IN_MAPS_ONCE]; - _FLOAT dat[MLO_N_LCL_IN_MAPS_ONCE]; - _FLOAT dat2[MLO_N_LCL_IN_MAPS_ONCE]; - -// - -// ATOMIC is needed if INPUTS in many waves -#if(MLO_N_LCL_IN_MAPS != MLO_N_INPUTS) - - if(in_grp_block == 0) - { - uint gbl_out_off = batch_id * MLO_OUT_BATCH_STRIDE + out_id * MLO_OUT_CHANNEL_STRIDE + pos; - __global _FLOAT* q = out_ptr + gbl_out_off; - - for(uint o = 0; o < MLO_N_LCL_OUT_MAPS; ++o) - { - *q = (_FLOAT)0; - q += MLO_OUT_CHANNEL_STRIDE; - } - } -#endif - - for(uint o = 0; o < MLO_N_LCL_OUT_MAPS; ++o) - { - accum[o] = (_FLOAT)0; - } - -#if MLO_N_INPUTS == ((MLO_N_INPUTS / MLO_N_LCL_IN_MAPS) * MLO_N_LCL_IN_MAPS) - // if(1) - int loops = MLO_CLOOP0; - -#if MLO_CHEAT_SHADER_COMPILER == 1 - // cheat shader compiler to disable loop unroll. it will have better SQC performance - if(grp_id2 == 0x1F) - { - loops = 377; // strange not to unroll loop - } -#endif -#else - int loops = MLO_CLOOP0; - - if(in_grp_block == (MLO_N_IN_GROUPS - 1)) - { - loops = MLO_CLOOP2; - } - -#if MLO_CHEAT_SHADER_COMPILER == 1 - // cheat shader compiler to disable loop unroll. it will have better SQC performance - if(grp_id2 == 0x1F) - { - loops = 377; // strange not to unroll loop - } -#endif - -#endif - { - __global const _FLOAT* p = in_ptr + gbl_in_off; - __constant _FLOAT* w = wei_ptr + wei_off; - - // read data - for(uint j = 0; j < MLO_N_LCL_IN_MAPS_ONCE; ++j) - { - - dat[j] = *p; - p += MLO_IN_CHANNEL_STRIDE; - } - - for(uint ci = 0; ci < (loops - 2); ci += 2) - { - // read data - for(uint j = 0; j < MLO_N_LCL_IN_MAPS_ONCE; ++j) - { - dat2[j] = *p; - p += MLO_IN_CHANNEL_STRIDE; - } - - // convolve - __constant _FLOAT* w1 = w; - for(uint o = 0; o < MLO_N_LCL_OUT_MAPS; ++o) - { - - __constant _FLOAT* w2 = w1; - - for(uint j = 0; j < MLO_N_LCL_IN_MAPS_ONCE; ++j) - { - weights[j] = *w2; - w2++; - } - w1 += MLO_WEI_CHANNEL_STRIDE; - - for(uint c = 0; c < MLO_N_LCL_IN_MAPS_ONCE; ++c) - { - accum[o] += dat[c] * weights[c]; - } - } - - // move weights offset - w += MLO_N_LCL_IN_MAPS_ONCE; - - // convolve - w1 = w; - for(uint j = 0; j < MLO_N_LCL_IN_MAPS_ONCE; ++j) - { - dat[j] = *p; - p += MLO_IN_CHANNEL_STRIDE; - } - - for(uint o = 0; o < MLO_N_LCL_OUT_MAPS; ++o) - { - - __constant _FLOAT* w2 = w1; - - for(uint j = 0; j < MLO_N_LCL_IN_MAPS_ONCE; ++j) - { - weights[j] = *w2; - w2++; - } - w1 += MLO_WEI_CHANNEL_STRIDE; - - for(uint c = 0; c < MLO_N_LCL_IN_MAPS_ONCE; ++c) - { - accum[o] += dat2[c] * weights[c]; - } - } - - // move weights offset - w += MLO_N_LCL_IN_MAPS_ONCE; - } - - // - // last 2 iterations - { // read data - for(uint j = 0; j < MLO_N_LCL_IN_MAPS_ONCE; ++j) - { - dat2[j] = *p; - p += MLO_IN_CHANNEL_STRIDE; - } - - // convolve - __constant _FLOAT* w1 = w; - for(uint o = 0; o < MLO_N_LCL_OUT_MAPS; ++o) - { - - __constant _FLOAT* w2 = w1; - - for(uint j = 0; j < MLO_N_LCL_IN_MAPS_ONCE; ++j) - { - weights[j] = *w2; - w2++; - } - w1 += MLO_WEI_CHANNEL_STRIDE; - - for(uint c = 0; c < MLO_N_LCL_IN_MAPS_ONCE; ++c) - { - accum[o] += dat[c] * weights[c]; - } - } - - // move weights offset - w += MLO_N_LCL_IN_MAPS_ONCE; - - // convolve - w1 = w; - - for(uint o = 0; o < MLO_N_LCL_OUT_MAPS; ++o) - { - - __constant _FLOAT* w2 = w1; - - for(uint j = 0; j < MLO_N_LCL_IN_MAPS_ONCE; ++j) - { - weights[j] = *w2; - w2++; - } - w1 += MLO_WEI_CHANNEL_STRIDE; - - for(uint c = 0; c < MLO_N_LCL_IN_MAPS_ONCE; ++c) - { - accum[o] += dat2[c] * weights[c]; - } - } - - // move weights offset - w += MLO_N_LCL_IN_MAPS_ONCE; - } - } - - uint gbl_out_off = batch_id * MLO_OUT_BATCH_STRIDE + out_id * MLO_OUT_CHANNEL_STRIDE + pos; - __global _FLOAT* q = out_ptr + gbl_out_off; - - for(uint o = 0; o < MLO_N_LCL_OUT_MAPS; ++o) - { -#if(MLO_N_LCL_IN_MAPS == MLO_N_INPUTS) - *q = accum[o]; - q += MLO_OUT_CHANNEL_STRIDE; -#else - AtomicAdd(q, accum[o]); - q += MLO_OUT_CHANNEL_STRIDE; -#endif - } -} diff --git a/projects/miopen/src/kernels/MIOpenConv1x1J1_stride.cl b/projects/miopen/src/kernels/MIOpenConv1x1J1_stride.cl deleted file mode 100644 index c2f97c7a5553..000000000000 --- a/projects/miopen/src/kernels/MIOpenConv1x1J1_stride.cl +++ /dev/null @@ -1,410 +0,0 @@ -/******************************************************************************* - * - * MIT License - * - * Copyright (c) 2017 Advanced Micro Devices, Inc. - * - * 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. - * - *******************************************************************************/ - -// Trying to use float ATOMIC_ADD to increase total waves -// For example, 7x7 : 49 -// Global 0: H * W * N * (K/ 16 outplane per threads ) * (( Inputlanes/ 128 ) or 1) -// N * H * W = -// [groupId / ((K/ 16 outplane per threads ) * (( Inputlanes/ 128 ) or 1))] * 64 + localId -// For example, H7*W7*N*C256*K64 -// FIRST 4 waves -// Hit L2 for read -// try to Hit L2 for Atomic_ADD -// Assembly Shader can utlize LDS to Reductiion -// However Shader Compiler of OpenCL will compile any localId[1] to FLAT_BUFFER_LOAD not constant -// load - -// WAVE 0 == N0_7x7 + N1_(7x2+1), output 0-15 from K=64, Inputplanes from 0-127 -// WAVE 1 == N0_7x7 + N1_(7x2+1), output 16-31 from K=64, Inputplanes from 0-127 -// WAVE 2 == N0_7x7 + N1_(7x2+1), output 31-47 from K=64, Inputplanes from 0-127 -// WAVE 3 == N0_7x7 + N1_(7x2+1), output 48-63 from K=64, Inputplanes from 0-127 -// 2nd 4 waves -// Hit L2 for read -// try to Hit L2 for Atomic_ADD - -// WAVE 4 == N0_7x7 + N1_(7x2+1), output 0-15 from K=64, Inputplanes from 128-255 -// WAVE 5 == N0_7x7 + N1_(7x2+1), output 16-31 from K=64, Inputplanes from 128-255 -// WAVE 6 == N0_7x7 + N1_(7x2+1), output 31-47 from K=64, Inputplanes from 128-255 -// WAVE 7 == N0_7x7 + N1_(7x2+1), output 48-63 from K=64, Inputplanes from 128-255 - -// STRIDE Mode: Global 0: H_out * W_out * N * (K/ 16 outplane per threads ) * (( Inputlanes/ 128 ) -// or 1) - -// example -#if 0 // ndef MLopen_RUNNING -#define MLO_FILTER_STRIDE0 2 -#define MLO_FILTER_STRIDE1 2 -#define MLO_N_LCL_IN_MAPS_ONCE 8 - -#define H 28 -#define W 28 -#define C 192 -#define K 64 - -#define MLO_IN_HEIGHT H -#define MLO_IN_WIDTH W -#define MLO_N_INPUTS C - -//128 or MLO_N_INPUTS -#define MLO_N_LCL_IN_MAPS 192 - -#define MLO_N_OUTPUTS K - -#define H_out 28 -#define W_out 28 -#define MLO_N_LCL_OUT_MAPS 16 - -#define MLO_N_IN_GROUPS ((MLO_N_INPUTS + MLO_N_LCL_IN_MAPS - 1) / MLO_N_LCL_IN_MAPS) -#define MLO_CLOOP0 (MLO_N_LCL_IN_MAPS / MLO_N_LCL_IN_MAPS_ONCE) -#define MLO_CLOOP2 \ - ((MLO_N_INPUTS - MLO_N_LCL_IN_MAPS * (MLO_N_IN_GROUPS - 1)) / MLO_N_LCL_IN_MAPS_ONCE) -#define MLO_CHEAT_SHADER_COMPILER 1 - -#endif - -#define MLO_IN_CHANNEL_STRIDE (H * W) -#define MLO_IN_BATCH_STRIDE (H * W * C) - -#define MLO_WEI_BSTRIDE (1 * 1 * C * K) -#define MLO_WEI_CHANNEL_STRIDE (1 * 1 * C) - -#define MLO_OUT_BATCH_STRIDE (H_out * W_out * K) -#define MLO_OUT_CHANNEL_STRIDE (H_out * W_out) - -#define FIXED_WORKGROUP_SIZE 64 - -#define MLO_N_OUT_GROUPS (MLO_N_OUTPUTS / MLO_N_LCL_OUT_MAPS) - -#define MLO_GRP_SZ0 64 -#define MLO_GRP_SZ1 1 -#define MLO_GRP_SZ2 1 - -#define PPCAT_NX(A, B) A##B -#define PPCAT(A, B) PPCAT_NX(A, B) -#define TWO 2 -#define FOUR 4 -#define EIGHT 8 - -#if MIOPEN_USE_FP16 == 1 -#pragma OPENCL EXTENSION cl_khr_fp16 : enable -#define _FLOAT half -#define _UNION_FLOAT_T half2 -#define INIT(A) ((half2)(A[0], A[1])) -#ifndef HALF_MAX -#define MAX_VAL 65504 /* max value */ -#else -#define MAX_VAL HALF_MAX -#endif -#endif -#if MIOPEN_USE_FP32 == 1 -#define _FLOAT float -#define _UNION_FLOAT_T float -#define INIT(A) (A[0]) -#ifndef FLT_MAX -#define MAX_VAL 3.402823466e+38F /* max value */ -#else -#define MAX_VAL FLT_MAX -#endif -#endif - -#define _FLOAT2 PPCAT(_FLOAT, TWO) -#define _FLOAT4 PPCAT(_FLOAT, FOUR) -#define _FLOAT8 PPCAT(_FLOAT, EIGHT) - -#define MLO_CONV_BIAS 0 -#define UNUSED __attribute((__unused__)) - -typedef union -{ - unsigned int intVal; - _UNION_FLOAT_T floatVal; -} starVal; - -inline void AtomicAdd(volatile __global _FLOAT* source, const _FLOAT operand) -{ - starVal newVal, prevVal; - - prevVal.floatVal = INIT(source); - while(true) - { -#if MIOPEN_USE_FP16 == 1 - newVal.floatVal = (_FLOAT2)(prevVal.floatVal.x + operand, source[1]); -#endif -#if MIOPEN_USE_FP32 == 1 - newVal.floatVal = prevVal.floatVal + operand; -#endif - newVal.intVal = - atomic_cmpxchg((volatile __global unsigned int*)source, prevVal.intVal, newVal.intVal); - - // equal to pass - if(newVal.intVal == prevVal.intVal) - break; - - prevVal.intVal = newVal.intVal; - } -} - -__attribute__((reqd_work_group_size(MLO_GRP_SZ0, MLO_GRP_SZ1, MLO_GRP_SZ2))) __kernel void -MIOpenConv1x1(const __global _FLOAT* __restrict in_ptr, - __constant _FLOAT* __restrict wei_ptr, -#if MLO_CONV_BIAS - const __global _FLOAT* __restrict bias, -#endif - __global _FLOAT* __restrict out_ptr, - UNUSED _FLOAT dummy_val // nothing -) -{ - - uint grp_id0 = get_group_id(0); - uint out_grp_block = grp_id0 % MLO_N_OUT_GROUPS; - uint in_grp_block = (uint)(grp_id0 / MLO_N_OUT_GROUPS) % MLO_N_IN_GROUPS; - uint grp_id0_faked = (uint)(grp_id0 / MLO_N_OUT_GROUPS) / MLO_N_IN_GROUPS; - - uint local_id0 = get_local_id(0); -#if MLO_CHEAT_SHADER_COMPILER == 1 - uint grp_id2 = get_group_id(2); -#endif - - uint pos = (grp_id0_faked * FIXED_WORKGROUP_SIZE + local_id0) % MLO_OUT_CHANNEL_STRIDE; - uint batch_id = (grp_id0_faked * FIXED_WORKGROUP_SIZE + local_id0) / MLO_OUT_CHANNEL_STRIDE; - - if(batch_id >= BATCHSIZE) - return; - - uint out_id = out_grp_block * MLO_N_LCL_OUT_MAPS; - - short out_pos_x = pos % W_out; - short out_pos_y = pos / W_out; - - uint in_pos = out_pos_x * MLO_FILTER_STRIDE0 + out_pos_y * MLO_FILTER_STRIDE1 * W; - - uint gbl_in_off = batch_id * MLO_IN_BATCH_STRIDE + - in_grp_block * MLO_N_LCL_IN_MAPS * MLO_IN_CHANNEL_STRIDE + in_pos; - - uint wei_off = out_id * MLO_WEI_CHANNEL_STRIDE + in_grp_block * MLO_N_LCL_IN_MAPS; - - _FLOAT accum[MLO_N_LCL_OUT_MAPS]; - _FLOAT weights[MLO_N_LCL_IN_MAPS_ONCE]; - _FLOAT dat[MLO_N_LCL_IN_MAPS_ONCE]; - _FLOAT dat2[MLO_N_LCL_IN_MAPS_ONCE]; - -// - -// ATOMIC is needed if INPUTS in many waves -#if(MLO_N_LCL_IN_MAPS != MLO_N_INPUTS) - - if(in_grp_block == 0) - { - uint gbl_out_off = batch_id * MLO_OUT_BATCH_STRIDE + out_id * MLO_OUT_CHANNEL_STRIDE + pos; - __global _FLOAT* q = out_ptr + gbl_out_off; - - for(uint o = 0; o < MLO_N_LCL_OUT_MAPS; ++o) - { - *q = (_FLOAT)0; - q += MLO_OUT_CHANNEL_STRIDE; - } - } -#endif - - for(uint o = 0; o < MLO_N_LCL_OUT_MAPS; ++o) - { - accum[o] = (_FLOAT)0; - } - -#if MLO_N_INPUTS == ((MLO_N_INPUTS / MLO_N_LCL_IN_MAPS) * MLO_N_LCL_IN_MAPS) - // if(1) - int loops = MLO_CLOOP0; - -#if MLO_CHEAT_SHADER_COMPILER == 1 - // cheat shader compiler to disable loop unroll. it will have better SQC performance - if(grp_id2 == 0x1F) - { - loops = 377; // strange not to unroll loop - } -#endif -#else - int loops = MLO_CLOOP0; - - if(in_grp_block == (MLO_N_IN_GROUPS - 1)) - { - loops = MLO_CLOOP2; - } - -#if MLO_CHEAT_SHADER_COMPILER == 1 - // cheat shader compiler to disable loop unroll. it will have better SQC performance - if(grp_id2 == 0x1F) - { - loops = 377; // strange not to unroll loop - } -#endif - -#endif - { - __global const _FLOAT* p = in_ptr + gbl_in_off; - __constant _FLOAT* w = wei_ptr + wei_off; - - // read data - for(uint j = 0; j < MLO_N_LCL_IN_MAPS_ONCE; ++j) - { - - dat[j] = *p; - p += MLO_IN_CHANNEL_STRIDE; - } - - for(uint ci = 0; ci < (loops - 2); ci += 2) - { - // read data - for(uint j = 0; j < MLO_N_LCL_IN_MAPS_ONCE; ++j) - { - dat2[j] = *p; - p += MLO_IN_CHANNEL_STRIDE; - } - - // convolve - __constant _FLOAT* w1 = w; - for(uint o = 0; o < MLO_N_LCL_OUT_MAPS; ++o) - { - - __constant _FLOAT* w2 = w1; - - for(uint j = 0; j < MLO_N_LCL_IN_MAPS_ONCE; ++j) - { - weights[j] = *w2; - w2++; - } - w1 += MLO_WEI_CHANNEL_STRIDE; - - for(uint c = 0; c < MLO_N_LCL_IN_MAPS_ONCE; ++c) - { - accum[o] += dat[c] * weights[c]; - } - } - - // move weights offset - w += MLO_N_LCL_IN_MAPS_ONCE; - - // convolve - w1 = w; - for(uint j = 0; j < MLO_N_LCL_IN_MAPS_ONCE; ++j) - { - dat[j] = *p; - p += MLO_IN_CHANNEL_STRIDE; - } - - for(uint o = 0; o < MLO_N_LCL_OUT_MAPS; ++o) - { - - __constant _FLOAT* w2 = w1; - - for(uint j = 0; j < MLO_N_LCL_IN_MAPS_ONCE; ++j) - { - weights[j] = *w2; - w2++; - } - w1 += MLO_WEI_CHANNEL_STRIDE; - - for(uint c = 0; c < MLO_N_LCL_IN_MAPS_ONCE; ++c) - { - accum[o] += dat2[c] * weights[c]; - } - } - - // move weights offset - w += MLO_N_LCL_IN_MAPS_ONCE; - } - - // - // last 2 iterations - { // read data - for(uint j = 0; j < MLO_N_LCL_IN_MAPS_ONCE; ++j) - { - dat2[j] = *p; - p += MLO_IN_CHANNEL_STRIDE; - } - - // convolve - __constant _FLOAT* w1 = w; - for(uint o = 0; o < MLO_N_LCL_OUT_MAPS; ++o) - { - - __constant _FLOAT* w2 = w1; - - for(uint j = 0; j < MLO_N_LCL_IN_MAPS_ONCE; ++j) - { - weights[j] = *w2; - w2++; - } - w1 += MLO_WEI_CHANNEL_STRIDE; - - for(uint c = 0; c < MLO_N_LCL_IN_MAPS_ONCE; ++c) - { - accum[o] += dat[c] * weights[c]; - } - } - - // move weights offset - w += MLO_N_LCL_IN_MAPS_ONCE; - - // convolve - w1 = w; - - for(uint o = 0; o < MLO_N_LCL_OUT_MAPS; ++o) - { - - __constant _FLOAT* w2 = w1; - - for(uint j = 0; j < MLO_N_LCL_IN_MAPS_ONCE; ++j) - { - weights[j] = *w2; - w2++; - } - w1 += MLO_WEI_CHANNEL_STRIDE; - - for(uint c = 0; c < MLO_N_LCL_IN_MAPS_ONCE; ++c) - { - accum[o] += dat2[c] * weights[c]; - } - } - - // move weights offset - w += MLO_N_LCL_IN_MAPS_ONCE; - } - } - - uint gbl_out_off = batch_id * MLO_OUT_BATCH_STRIDE + out_id * MLO_OUT_CHANNEL_STRIDE + pos; - __global _FLOAT* q = out_ptr + gbl_out_off; - - for(uint o = 0; o < MLO_N_LCL_OUT_MAPS; ++o) - { -#if(MLO_N_LCL_IN_MAPS == MLO_N_INPUTS) - *q = accum[o]; - q += MLO_OUT_CHANNEL_STRIDE; -#else - AtomicAdd(q, accum[o]); - q += MLO_OUT_CHANNEL_STRIDE; -#endif - } -} diff --git a/projects/miopen/src/kernels/MIOpenConv1x1S.cl b/projects/miopen/src/kernels/MIOpenConv1x1S.cl deleted file mode 100644 index a4d9d937e977..000000000000 --- a/projects/miopen/src/kernels/MIOpenConv1x1S.cl +++ /dev/null @@ -1,408 +0,0 @@ -/******************************************************************************* - * - * MIT License - * - * Copyright (c) 2017 Advanced Micro Devices, Inc. - * - * 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. - * - *******************************************************************************/ - -#include "float_types.h" -#include "math_ops.h" - -// Since float_types.h has enabled true mixed precision for all -// direct ocl kernels, this kernel needs to retain its older behavior as it is -// dependent upon tunability which isn't slated for MIOpen 2.0 PR #1725 -#if MIOPEN_USE_FP16 == 1 -#define _FLOAT_PREC _FLOAT -#else -#define _FLOAT_PREC _FLOAT_ACCUM -#endif - -#define UNUSED __attribute__((__unused__)) - -#define DBG_OUT_OF_RNGE 0 - -__attribute__((reqd_work_group_size(MLO_GRP_SZ0, MLO_GRP_SZ1, MLO_GRP_SZ2))) __kernel void -MIOpenConv1x1(const __global _FLOAT* __restrict in_ptr, - __constant _FLOAT* __restrict wei_ptr, -#if MLO_CONV_BIAS - const __global _FLOAT* __restrict bias, -#endif - __global _FLOAT* __restrict out_ptr, - UNUSED _FLOAT dummy_val // nothing -) -{ - _FLOAT weights[MLO_N_LCL_OUT_MAPS][MLO_N_LCL_IN_MAPS]; - - uint gbl_id0 = get_global_id(0); - uint batch_id = gbl_id0 / MLO_MAP_SZ4; // batch - uint pos = gbl_id0 % MLO_MAP_SZ4; - - uint out_grp_block = get_group_id(1); // block of outputs for the entire group - uint out_id = out_grp_block * MLO_N_LCL_OUT_MAPS; - - uint gbl_in_off = batch_id * MLO_IN_BATCH_STRIDE + pos * MLO_READ_UNIT; - - uint wei_off = out_id * -#if MLO_DIR_FORWARD == 1 - MLO_WEI_BSTRIDE -#else - MLO_WEI_CHANNEL_STRIDE -#endif - ; - - _FLOAT_PREC accum[MLO_N_LCL_OUT_MAPS][MLO_READ_UNIT] = {{(_FLOAT_PREC)0.0f}}; - _FLOAT dat[MLO_N_LCL_IN_MAPS][MLO_READ_UNIT]; - - for(uint o = 0; o < MLO_N_LCL_OUT_MAPS; ++o) - { - for(uint i = 0; i < MLO_READ_UNIT; ++i) - { - accum[o][i] = (_FLOAT_PREC)0.0f; - } - } - - for(uint ci = 0; ci < MLO_CLOOP0; ++ci, - // move input offset - gbl_in_off += MLO_N_LCL_IN_MAPS * MLO_IN_CHANNEL_STRIDE, - - // move weights offset - wei_off += MLO_N_LCL_IN_MAPS * -#if MLO_DIR_FORWARD == 1 - MLO_WEI_CHANNEL_STRIDE -#else - MLO_WEI_BSTRIDE -#endif - ) - { - // read weights - - __constant _FLOAT* w1 = wei_ptr + wei_off; - - for(uint o = 0; o < MLO_N_LCL_OUT_MAPS; ++o, - w1 += -#if MLO_DIR_FORWARD == 1 - MLO_WEI_BSTRIDE -#else - MLO_WEI_CHANNEL_STRIDE -#endif - ) - { - __constant _FLOAT* w2 = w1; - for(uint c = 0; c < MLO_N_LCL_IN_MAPS; ++c, - w2 += -#if MLO_DIR_FORWARD == 1 - MLO_WEI_CHANNEL_STRIDE -#else - MLO_WEI_BSTRIDE -#endif - ) - { - - weights[o][c] = *w2; - -#if DBG_OUT_OF_RNGE - if(wei_off2 >= MLO_N_INPUTS * MLO_N_OUTPUTS) - { - printf("K:oor: weights\n"); - } -#endif - } - } - - // convolve with all weights - // read data - // Shader compiler will use GLOAL_LOAD_DWORD's OFFSET for *(ptr+index) access - // Shader compiler will not use GLOAL_LOAD_DWORD's OFFSET for ptr[index] access - - __global const _FLOAT* p = in_ptr + gbl_in_off; - - for(uint j = 0; j < MLO_N_LCL_IN_MAPS; j++) - { - for(uint i = 0; i < MLO_READ_UNIT; ++i) - { - dat[j][i] = *(p + i); -#if DBG_OUT_OF_RNGE - if(gbl_in_off1 + i >= MLO_IN_BATCH_STRIDE * MLO_BATCH_SZ) - { - printf("K:oor: inputs\n"); - } -#endif - } - p += MLO_IN_CHANNEL_STRIDE; - } - - // convolve - for(uint o = 0; o < MLO_N_LCL_OUT_MAPS; ++o) - { - _FLOAT_PREC acc[MLO_READ_UNIT] = {(_FLOAT_PREC)0.0f}; - for(uint c = 0; c < MLO_N_LCL_IN_MAPS; ++c) - { - _FLOAT_PREC we = CVT_FLOAT2ACCUM(weights[o][c]); - _FLOAT* d = &dat[c][0]; - for(uint i = 0; i < MLO_READ_UNIT; ++i) - { - acc[i] += CVT_FLOAT2ACCUM(d[i]) * we; - } - } - for(uint i = 0; i < MLO_READ_UNIT; ++i) - accum[o][i] += acc[i]; - } - } - - uint gbl_out_off = - batch_id * MLO_OUT_BATCH_STRIDE + pos * MLO_READ_UNIT + out_id * MLO_OUT_CHANNEL_STRIDE; - __global _FLOAT* q = out_ptr + gbl_out_off; - - for(uint o = 0; o < MLO_N_LCL_OUT_MAPS; ++o, q += MLO_OUT_CHANNEL_STRIDE) - { - for(uint i = 0; i < MLO_READ_UNIT; ++i) - { - *(q + i) = CVT_ACCUM2FLOAT(accum[o][i]); - } - } -} - -/************************************************************************ -stride and padding -*************************************************************************/ -__attribute__((reqd_work_group_size(MLO_GRP_SZ0, MLO_GRP_SZ1, MLO_GRP_SZ2))) __kernel void -MIOpenConv1x1pquv(const __global _FLOAT* __restrict in_ptr, - __constant _FLOAT* __restrict wei_ptr, -#if MLO_CONV_BIAS - const __global _FLOAT* __restrict bias, -#endif - __global _FLOAT* __restrict out_ptr, - UNUSED _FLOAT dummy_val // nothing -) -{ - - _FLOAT weights[MLO_N_LCL_OUT_MAPS][MLO_N_LCL_IN_MAPS]; - - uint gbl_id0 = get_global_id(0); - - uint batch_id = gbl_id0 / MLO_MAP_SZ4; // batch - uint pos = gbl_id0 % MLO_MAP_SZ4; - uint pos_out_y = pos / MLO_OUT_WIDTH4; - uint pos_out_x = pos % MLO_OUT_WIDTH4; - -#if MLO_DIR_FORWARD == 1 - uint pos_in_y = pos_out_y * MLO_FILTER_STRIDE1; - uint pos_in_x = pos_out_x * MLO_FILTER_STRIDE0; -#else - uint pos_in_y = pos_out_y; /// MLO_FILTER_STRIDE1; - divided already - uint pos_in_x = pos_out_x; // MLO_FILTER_STRIDE0; - divided already -#endif - - uint out_grp_block = get_group_id(1); // block of outputs for the entire group - uint out_id = out_grp_block * MLO_N_LCL_OUT_MAPS; - - uint gbl_in_off = - batch_id * MLO_IN_BATCH_STRIDE + pos_in_y * MLO_IN_STRIDE + pos_in_x * MLO_READ_UNIT; - // bool vis = (pos_in_y < MLO_IN_HEIGHT); - // gbl_in_off = (vis) ? gbl_in_off : 0; - - uint wei_off = out_id * -#if MLO_DIR_FORWARD == 1 - MLO_WEI_BSTRIDE -#else - MLO_WEI_CHANNEL_STRIDE -#endif - ; - - _FLOAT_PREC accum[MLO_N_LCL_OUT_MAPS][MLO_READ_UNIT]; - _FLOAT dat[MLO_N_LCL_IN_MAPS][MLO_READ_UNIT]; - - for(uint o = 0; o < MLO_N_LCL_OUT_MAPS; ++o) - { - for(uint i = 0; i < MLO_READ_UNIT; ++i) - { - accum[o][i] = CVT_FLOAT2ACCUM(0); - } - } - - const __global _FLOAT* i_ptr = in_ptr + gbl_in_off; - __constant _FLOAT* w_ptr = wei_ptr + wei_off; - for(uint ci = 0; ci < MLO_CLOOP0; ++ci) - { - - // convolve with all weights - // read data - - for(uint j = 0; j < MLO_N_LCL_IN_MAPS; ++j) - { - uint i = 0; -#if MLO_READ_UNIT > 1 - for(; i < MLO_READ_UNIT - 1; ++i) - { - uint off = i -#if MLO_DIR_FORWARD == 1 - * MLO_FILTER_STRIDE0 -#endif - ; - dat[j][i] = *(i_ptr + off); - } -#endif - - for(; i < MLO_READ_UNIT; ++i) - { - // vis &= (pos_in_x + i*MLO_FILTER_STRIDE0 < - // MLO_IN_WIDTH); - uint off = i -#if MLO_DIR_FORWARD == 1 - * MLO_FILTER_STRIDE0 -#endif - ; - // off = (vis) ? off : 0; - _FLOAT val = *(i_ptr + off); - dat[j][i] = val; - // dat[j][i] = (vis)? dat[j][i] : (_FLOAT)(0); - } - - i_ptr += MLO_IN_CHANNEL_STRIDE; - } - // read weights - __constant _FLOAT* w_ptr0 = w_ptr; - - for(uint o = 0; o < MLO_N_LCL_OUT_MAPS; ++o) - { - - __constant _FLOAT* w_ptr1 = w_ptr0; - - for(uint c = 0; c < MLO_N_LCL_IN_MAPS; ++c) - { - weights[o][c] = *w_ptr1; - w_ptr1 += -#if MLO_DIR_FORWARD == 1 - MLO_WEI_CHANNEL_STRIDE -#else - MLO_WEI_BSTRIDE -#endif - ; - } - - w_ptr0 += -#if MLO_DIR_FORWARD == 1 - MLO_WEI_BSTRIDE -#else - MLO_WEI_CHANNEL_STRIDE -#endif - ; - } - - w_ptr += MLO_N_LCL_IN_MAPS * -#if MLO_DIR_FORWARD == 1 - MLO_WEI_CHANNEL_STRIDE -#else - MLO_WEI_BSTRIDE -#endif - ; - // convolve - for(uint o = 0; o < MLO_N_LCL_OUT_MAPS; ++o) - { - for(uint c = 0; c < MLO_N_LCL_IN_MAPS; ++c) - { - for(uint i = 0; i < MLO_READ_UNIT; ++i) - { - accum[o][i] += CVT_FLOAT2ACCUM(dat[c][i]) * CVT_FLOAT2ACCUM(weights[o][c]); -#if 0 - if (pos_out_y == 2 && pos_out_x == 0) - { - printf((__constant char *)"K:c: %f %f %f %f\n", - accum[o][i], - dat[c][i] * weights[o][c], - dat[c][i], - weights[o][c] - ); - } -#endif - } - } - } - } - - uint out_y = pos_out_y -#if MLO_DIR_FORWARD == 0 - * MLO_FILTER_STRIDE1 -#endif - ; - uint out_x = pos_out_x -#if MLO_DIR_FORWARD == 0 - * MLO_FILTER_STRIDE0 -#endif - ; - - uint gbl_out_off = batch_id * MLO_OUT_BATCH_STRIDE + out_id * MLO_OUT_CHANNEL_STRIDE + - out_y * MLO_OUT_STRIDE + out_x * MLO_READ_UNIT; - - __global _FLOAT* q = out_ptr + gbl_out_off; - - for(uint o = 0; o < MLO_N_LCL_OUT_MAPS; ++o, q += MLO_OUT_CHANNEL_STRIDE) - { - - for(uint i = 0; i < MLO_READ_UNIT; ++i) - { - __global _FLOAT* q1 = q; - q1 += i -#if MLO_DIR_FORWARD == 0 - - * MLO_FILTER_STRIDE0 -#endif - ; - *q1 = CVT_ACCUM2FLOAT(accum[o][i]); - -#if MLO_DIR_FORWARD == 0 - for(uint s = 1; s < MLO_FILTER_STRIDE0; ++s) - { -#if MLO_HORIZ_ALIGNED == 0 - if(out_x + s < MLO_OUT_WIDTH) -#endif - { - *(q1 + s) = CVT_ACCUM2FLOAT(0); - } - } -#endif - } - -#if MLO_DIR_FORWARD == 0 - __global _FLOAT* q2 = q; - for(uint j = 1; j < MLO_FILTER_STRIDE1; ++j) - { - q2 += MLO_OUT_STRIDE; -#if MLO_VERT_ALIGNED == 0 - if(out_y + j < MLO_OUT_HEIGHT) -#endif - { - - for(uint s = 0; s < MLO_READ_UNIT * MLO_FILTER_STRIDE0; ++s) - { -#if MLO_HORIZ_ALIGNED == 0 - if(out_x + s < MLO_OUT_WIDTH) -#endif - { - *(q2 + s) = CVT_ACCUM2FLOAT(0); - } - } - } - } -#endif - } -} diff --git a/projects/miopen/src/kernels/gfx803_36.HIP.fdb.txt.bz2 b/projects/miopen/src/kernels/gfx803_36.HIP.fdb.txt.bz2 index 840ed3d826cf..732ea561702b 100644 Binary files a/projects/miopen/src/kernels/gfx803_36.HIP.fdb.txt.bz2 and b/projects/miopen/src/kernels/gfx803_36.HIP.fdb.txt.bz2 differ diff --git a/projects/miopen/src/kernels/gfx803_36.OpenCL.fdb.txt.bz2 b/projects/miopen/src/kernels/gfx803_36.OpenCL.fdb.txt.bz2 index fb32547b7b62..583fb32ffbe9 100644 Binary files a/projects/miopen/src/kernels/gfx803_36.OpenCL.fdb.txt.bz2 and b/projects/miopen/src/kernels/gfx803_36.OpenCL.fdb.txt.bz2 differ diff --git a/projects/miopen/src/kernels/gfx803_64.HIP.fdb.txt.bz2 b/projects/miopen/src/kernels/gfx803_64.HIP.fdb.txt.bz2 index a874d8a32df9..68526324537f 100644 Binary files a/projects/miopen/src/kernels/gfx803_64.HIP.fdb.txt.bz2 and b/projects/miopen/src/kernels/gfx803_64.HIP.fdb.txt.bz2 differ diff --git a/projects/miopen/src/kernels/gfx803_64.OpenCL.fdb.txt.bz2 b/projects/miopen/src/kernels/gfx803_64.OpenCL.fdb.txt.bz2 index fb32547b7b62..583fb32ffbe9 100644 Binary files a/projects/miopen/src/kernels/gfx803_64.OpenCL.fdb.txt.bz2 and b/projects/miopen/src/kernels/gfx803_64.OpenCL.fdb.txt.bz2 differ diff --git a/projects/miopen/src/kernels/gfx900_56.HIP.fdb.txt.bz2 b/projects/miopen/src/kernels/gfx900_56.HIP.fdb.txt.bz2 index 80da6940a6ed..9b395bcc6197 100644 Binary files a/projects/miopen/src/kernels/gfx900_56.HIP.fdb.txt.bz2 and b/projects/miopen/src/kernels/gfx900_56.HIP.fdb.txt.bz2 differ diff --git a/projects/miopen/src/kernels/gfx900_56.OpenCL.fdb.txt.bz2 b/projects/miopen/src/kernels/gfx900_56.OpenCL.fdb.txt.bz2 index 434af48becfd..55f18c4952ce 100644 Binary files a/projects/miopen/src/kernels/gfx900_56.OpenCL.fdb.txt.bz2 and b/projects/miopen/src/kernels/gfx900_56.OpenCL.fdb.txt.bz2 differ diff --git a/projects/miopen/src/kernels/gfx900_64.HIP.fdb.txt.bz2 b/projects/miopen/src/kernels/gfx900_64.HIP.fdb.txt.bz2 index a874d8a32df9..68526324537f 100644 Binary files a/projects/miopen/src/kernels/gfx900_64.HIP.fdb.txt.bz2 and b/projects/miopen/src/kernels/gfx900_64.HIP.fdb.txt.bz2 differ diff --git a/projects/miopen/src/kernels/gfx900_64.OpenCL.fdb.txt.bz2 b/projects/miopen/src/kernels/gfx900_64.OpenCL.fdb.txt.bz2 index 805399ae0aad..088561ae542e 100644 Binary files a/projects/miopen/src/kernels/gfx900_64.OpenCL.fdb.txt.bz2 and b/projects/miopen/src/kernels/gfx900_64.OpenCL.fdb.txt.bz2 differ diff --git a/projects/miopen/src/kernels/gfx906_60.HIP.fdb.txt.bz2 b/projects/miopen/src/kernels/gfx906_60.HIP.fdb.txt.bz2 index 33b4440ac3b5..e50263c4a089 100644 Binary files a/projects/miopen/src/kernels/gfx906_60.HIP.fdb.txt.bz2 and b/projects/miopen/src/kernels/gfx906_60.HIP.fdb.txt.bz2 differ diff --git a/projects/miopen/src/kernels/gfx906_60.OpenCL.fdb.txt.bz2 b/projects/miopen/src/kernels/gfx906_60.OpenCL.fdb.txt.bz2 index 9b7fef1ef201..86aad6add246 100644 Binary files a/projects/miopen/src/kernels/gfx906_60.OpenCL.fdb.txt.bz2 and b/projects/miopen/src/kernels/gfx906_60.OpenCL.fdb.txt.bz2 differ diff --git a/projects/miopen/src/kernels/gfx906_64.HIP.fdb.txt.bz2 b/projects/miopen/src/kernels/gfx906_64.HIP.fdb.txt.bz2 index 0debf71bf1c0..9ca4b4c4d69d 100644 Binary files a/projects/miopen/src/kernels/gfx906_64.HIP.fdb.txt.bz2 and b/projects/miopen/src/kernels/gfx906_64.HIP.fdb.txt.bz2 differ diff --git a/projects/miopen/src/kernels/gfx906_64.OpenCL.fdb.txt.bz2 b/projects/miopen/src/kernels/gfx906_64.OpenCL.fdb.txt.bz2 index 7b62d0f7f8ba..8749a967d7a6 100644 Binary files a/projects/miopen/src/kernels/gfx906_64.OpenCL.fdb.txt.bz2 and b/projects/miopen/src/kernels/gfx906_64.OpenCL.fdb.txt.bz2 differ diff --git a/projects/miopen/src/kernels/gfx90878.HIP.fdb.txt.bz2 b/projects/miopen/src/kernels/gfx90878.HIP.fdb.txt.bz2 index 9ac2d96f91df..3ed2561a2710 100644 Binary files a/projects/miopen/src/kernels/gfx90878.HIP.fdb.txt.bz2 and b/projects/miopen/src/kernels/gfx90878.HIP.fdb.txt.bz2 differ diff --git a/projects/miopen/src/kernels/gfx90a68.HIP.fdb.txt.bz2 b/projects/miopen/src/kernels/gfx90a68.HIP.fdb.txt.bz2 index 1938be9d8f2e..7397d60f6c60 100644 Binary files a/projects/miopen/src/kernels/gfx90a68.HIP.fdb.txt.bz2 and b/projects/miopen/src/kernels/gfx90a68.HIP.fdb.txt.bz2 differ diff --git a/projects/miopen/src/kernels/gfx90a6e.HIP.fdb.txt.bz2 b/projects/miopen/src/kernels/gfx90a6e.HIP.fdb.txt.bz2 index 961a41281b6a..23d54c558511 100644 Binary files a/projects/miopen/src/kernels/gfx90a6e.HIP.fdb.txt.bz2 and b/projects/miopen/src/kernels/gfx90a6e.HIP.fdb.txt.bz2 differ diff --git a/projects/miopen/src/mlo_dir_conv.cpp b/projects/miopen/src/mlo_dir_conv.cpp index 429dce86ac0d..b7ae7711f8d5 100644 --- a/projects/miopen/src/mlo_dir_conv.cpp +++ b/projects/miopen/src/mlo_dir_conv.cpp @@ -1,28 +1,5 @@ -/******************************************************************************* - * - * MIT License - * - * Copyright (c) 2026 Advanced Micro Devices, Inc. - * - * 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. - * - *******************************************************************************/ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT #include #include @@ -79,7 +56,6 @@ static auto GetDirectSolvers() miopen::solver::conv::ConvAsm5x10u2v2b1, miopen::solver::conv::ConvHipDirectFwd11x11, miopen::solver::conv::ConvOclDirectFwdGen, - miopen::solver::conv::ConvOclDirectFwd1x1, miopen::solver::conv::ConvOclDirectFwd, miopen::solver::conv::ConvDirectNaiveConvFwd, miopen::solver::conv::ConvDirectNaiveConvBwd, diff --git a/projects/miopen/src/solver.cpp b/projects/miopen/src/solver.cpp index bf2c3b82cad4..f4aca3cd353a 100644 --- a/projects/miopen/src/solver.cpp +++ b/projects/miopen/src/solver.cpp @@ -1,4 +1,4 @@ -// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// Copyright © Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT #include @@ -310,7 +310,7 @@ inline SolverRegistrar::SolverRegistrar(IdRegistryData& registry) Primitive::Fusion, fusion::ConvOclDirectFwdFused{}.SolverDbId(), miopenConvolutionAlgoDirect); - RegisterWithSolver(registry, ++id, conv::ConvOclDirectFwd1x1{}, miopenConvolutionAlgoDirect); + ++id; // removed ConvOclDirectFwd1x1 RegisterWithSolver(registry, ++id, conv::ConvBinWinograd3x3U{}, miopenConvolutionAlgoWinograd); RegisterWithSolver(registry, ++id, conv::ConvBinWinogradRxS{}, miopenConvolutionAlgoWinograd); RegisterWithSolver(registry, ++id, conv::ConvAsmBwdWrW3x3{}, miopenConvolutionAlgoDirect); diff --git a/projects/miopen/src/solver/conv/conv_ocl_dir2Dfwd1x1.cpp b/projects/miopen/src/solver/conv/conv_ocl_dir2Dfwd1x1.cpp deleted file mode 100644 index 4a73f6b8d2a3..000000000000 --- a/projects/miopen/src/solver/conv/conv_ocl_dir2Dfwd1x1.cpp +++ /dev/null @@ -1,433 +0,0 @@ -/******************************************************************************* - * - * MIT License - * - * Copyright (c) 2017 Advanced Micro Devices, Inc. - * - * 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. - * - *******************************************************************************/ - -#include -#include -#include -#include -#include -#include - -/// WORKAROUND_SWDEV_271887 disables ConvOclDirectFwd1x1 solver on gfx10 and gfx11 due to precision -/// issues. -#define WORKAROUND_SWDEV_271887 1 - -MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1) - -namespace miopen { -namespace solver { -namespace conv { - -using ProblemDescription = miopen::conv::ProblemDescription; - -bool ConvOclDirectFwd1x1::IsApplicable(const ExecutionContext& ctx, - const ProblemDescription& problem) const -{ -#if WORKAROUND_SWDEV_271887 - { - if(StartsWith(ctx.GetStream().GetDeviceName(), "gfx10") || - StartsWith(ctx.GetStream().GetDeviceName(), "gfx11")) - { - if(!env::enabled(MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1)) - return false; - } - } -#endif - const std::string name = ctx.GetStream().GetDeviceName(); - if(!(StartsWith(name, "gfx8") || StartsWith(name, "gfx90") || StartsWith(name, "gfx103"))) - return false; - if(env::disabled(MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1)) - return false; - if(!ctx.use_opencl_convolutions) - return false; - if(!problem.Is2d()) - return false; - if(!(problem.IsDirectionForward() || problem.IsDirectionBackwardData())) - return false; - if(problem.HasNonPackedTensors()) - return false; - if(!problem.AllTensorsDimsFitIntoInt()) - return false; - if(problem.IsAsymmetricPadH() || problem.IsAsymmetricPadW()) - return false; - if(!(problem.IsFp32() || problem.IsFp16() || problem.IsBfp16())) - return false; - if(problem.IsTensorsCasted()) - return false; - if(!problem.IsLayoutDefault()) - return false; - - return problem.GetDilationW() == 1 && problem.GetDilationH() == 1 && - problem.GetWeightsWidth() == 1 && problem.GetWeightsHeight() == 1 && - problem.GetGroupCount() == 1 && - // TODO: update 1x1 fwd kernel to support padding - problem.GetPadW() == 0 && problem.GetPadH() == 0; -} - -ConvSolution ConvOclDirectFwd1x1::GetSolution(const ExecutionContext& ctx, - const ProblemDescription& problem, - const LegacyPerformanceConfig& config) const -{ - ConvSolution result; - config.CopyTo(result); - - // if(problem.GetOutChannels() % 4 == 0 && problem.GetInChannels() % 4 == 0) - { - // int version = result.out_pix_tile1; - - if((problem.IsDirectionForward() && problem.GetInChannels() % 16 == 0 && - problem.GetOutChannels() % 16 == 0) && - (problem.GetInDataType() == miopenFloat)) - { - - int N_LCL_IN_MAPS = result.n_in_data_tiles; - - int N_LCL_OUT_MAPS = result.n_out_pix_tiles; - // 0 or 1 - int CHEAT_SHADER_COMPILER = result.out_pix_tile0; - - int BATCHSIZE = problem.GetBatchSize(); - int W = problem.GetInWidth(); - int H = problem.GetInHeight(); - int C = problem.GetInChannels(); - int K = problem.GetOutChannels(); - int W_out = problem.GetOutWidth(); - int H_out = problem.GetOutHeight(); - - N_LCL_IN_MAPS = std::min(N_LCL_IN_MAPS, C); - N_LCL_OUT_MAPS = std::min(N_LCL_OUT_MAPS, K); - - while((K % N_LCL_OUT_MAPS) != 0 && N_LCL_OUT_MAPS > 16) - { - N_LCL_OUT_MAPS /= 2; - } - - result.n_out_pix_tiles = N_LCL_OUT_MAPS; - - if(N_LCL_IN_MAPS < C && N_LCL_IN_MAPS > 0 && (N_LCL_IN_MAPS % 8) == 0) - { - // Pass will do nothing - } - else - { - N_LCL_IN_MAPS = C; - } - - result.n_in_data_tiles = N_LCL_IN_MAPS; - - /* - #define H 28 - #define W 28 - #define C 192 - #define K 64 - - #define MLO_IN_HEIGHT H - #define MLO_IN_WIDTH W - #define MLO_N_INPUTS C - - //128 or MLO_N_INPUTS - #define MLO_N_LCL_IN_MAPS 128 - - #define MLO_N_OUTPUTS K - - #define H_out H - #define W_out W - */ - // #define MLO_N_IN_GROUPS (( MLO_N_INPUTS + MLO_N_LCL_IN_MAPS - 1) / - // MLO_N_LCL_IN_MAPS) - // #define MLO_CLOOP0 (MLO_N_LCL_IN_MAPS/MLO_N_LCL_IN_MAPS_ONCE ) - // #define MLO_CLOOP2 ((MLO_N_INPUTS - - // MLO_N_LCL_IN_MAPS*(MLO_N_IN_GROUPS-1)) / MLO_N_LCL_IN_MAPS_ONCE ) - // #define MLO_N_LCL_OUT_MAPS 16 - - int N_IN_GROUPS = (C + N_LCL_IN_MAPS - 1) / N_LCL_IN_MAPS; - int N_LCL_IN_MAPS_ONCE = 8; - - if(problem.GetKernelStrideW() > 1 || problem.GetKernelStrideH() > 1) - N_LCL_IN_MAPS_ONCE = 4; - - int CLOOP0 = N_LCL_IN_MAPS / N_LCL_IN_MAPS_ONCE; - int CLOOP2 = (C - N_LCL_IN_MAPS * (N_IN_GROUPS - 1)) / N_LCL_IN_MAPS_ONCE; - - KernelInfo kernel; - - kernel.comp_options = - std::string(" -DMLO_N_LCL_IN_MAPS_ONCE=") + std::to_string(N_LCL_IN_MAPS_ONCE) + - std::string(" -DBATCHSIZE=") + std::to_string(BATCHSIZE) + std::string(" -DH=") + - std::to_string(H) + std::string(" -DW=") + std::to_string(W) + - std::string(" -DC=") + std::to_string(C) + std::string(" -DK=") + - std::to_string(K) + std::string(" -DMLO_N_LCL_IN_MAPS=") + - std::to_string(N_LCL_IN_MAPS) + std::string(" -DMLO_N_INPUTS=") + - std::to_string(C) + std::string(" -DMLO_N_OUTPUTS=") + std::to_string(K) + - std::string(" -DH_out=") + std::to_string(H_out) + std::string(" -DW_out=") + - std::to_string(W_out) + std::string(" -DMLO_N_IN_GROUPS=") + - std::to_string(N_IN_GROUPS) + std::string(" -DMLO_CLOOP0=") + - std::to_string(CLOOP0) + std::string(" -DMLO_CLOOP2=") + std::to_string(CLOOP2) + - std::string(" -DMLO_N_LCL_OUT_MAPS=") + std::to_string(N_LCL_OUT_MAPS) + - std::string(" -DMLO_CHEAT_SHADER_COMPILER=") + - std::to_string(CHEAT_SHADER_COMPILER) + - std::string( - " -DMLopen_RUNNING=1") + // to disable macro defines for CodeXL Shader Analyzer - ctx.general_compile_options; - - kernel.comp_options = std::string(" -DMLO_FILTER_STRIDE0=") + - std::to_string(problem.GetKernelStrideW()) + - std::string(" -DMLO_FILTER_STRIDE1=") + - std::to_string(problem.GetKernelStrideH()) + kernel.comp_options; - - // std::cout << "compile options:\n"<< _comp_options << std::endl; - - // 1x1_Stride: FIX ME!!! NO padding support - if(problem.GetKernelStrideW() > 1 || problem.GetKernelStrideH() > 1) - { - int FIXED_WORKGROUP_SIZE = 64; - - size_t N_OUT_GROUPS = (K / N_LCL_OUT_MAPS); - - size_t local_wk1 = 1; - kernel.l_wk.push_back(FIXED_WORKGROUP_SIZE); - kernel.l_wk.push_back(local_wk1); - kernel.l_wk.push_back(1); - - size_t imagesizeAlign = - ((problem.GetOutWidth() * problem.GetOutHeight() * problem.GetBatchSize() + - FIXED_WORKGROUP_SIZE - 1) / - FIXED_WORKGROUP_SIZE) * - FIXED_WORKGROUP_SIZE; - - size_t gbl_wk0 = imagesizeAlign * N_IN_GROUPS * N_OUT_GROUPS; - size_t gbl_wk1 = local_wk1; - size_t gbl_wk2 = 1; - - kernel.g_wk.push_back(gbl_wk0); - kernel.g_wk.push_back(gbl_wk1); - kernel.g_wk.push_back(gbl_wk2); - - kernel.kernel_file = "MIOpenConv1x1J1_stride.cl"; - kernel.kernel_name = "MIOpenConv1x1"; - result.construction_params.push_back(kernel); - } - else - { - int FIXED_WORKGROUP_SIZE = 64; - - kernel.l_wk.push_back(FIXED_WORKGROUP_SIZE); - kernel.l_wk.push_back(1); - kernel.l_wk.push_back(1); - - size_t imagesizeAlign = - ((problem.GetInWidth() * problem.GetInHeight() * problem.GetBatchSize() + - FIXED_WORKGROUP_SIZE - 1) / - FIXED_WORKGROUP_SIZE) * - FIXED_WORKGROUP_SIZE; - size_t N_OUT_GROUPS = (K / N_LCL_OUT_MAPS); - - size_t gbl_wk0 = imagesizeAlign * N_IN_GROUPS * N_OUT_GROUPS; - - size_t gbl_wk1 = 1; - - size_t gbl_wk2 = 1; - - kernel.g_wk.push_back(gbl_wk0); - kernel.g_wk.push_back(gbl_wk1); - kernel.g_wk.push_back(gbl_wk2); - - kernel.kernel_file = "MIOpenConv1x1J1.cl"; - kernel.kernel_name = "MIOpenConv1x1"; - result.construction_params.push_back(kernel); - } - // std::cout << _kernel_file << std::endl; - } - else - { - - // parameters - // int i_sz = problem.GetInWidth() * problem.GetInHeight(); - // _out_pix_tile0 = (i_sz & 1) ? 1 : 2; - result.out_pix_tile0 = - std::min(static_cast(problem.GetOutWidth()), result.out_pix_tile0); - result.out_pix_tile1 = - std::min(static_cast(problem.GetOutHeight()), result.out_pix_tile1); - if(!problem.IsDirectionForward()) - { - while(problem.GetOutWidth() % result.out_pix_tile0 != 0 && result.out_pix_tile0 > 1) - { - result.out_pix_tile0 /= 2; - } - } - - int read_unit = result.out_pix_tile0; - while(problem.GetInWidth() % read_unit != 0 && read_unit > 1) - { - read_unit /= 2; - } - - // problem.GetOutWidth() - // _n_out_pix_tiles = 16; - // _n_in_data_tiles = 4; - // _grp_tile0 = 64; - - int wei_cstride = problem.GetWeightsWidth() * problem.GetWeightsHeight(); - // backward: inputs are forward outputs - const bool is_forward = problem.IsDirectionForward(); - int wei_bstride = - (is_forward ? problem.GetInChannels() : problem.GetOutChannels()) * wei_cstride; - - int OUT_WIDTH4 = problem.GetOutWidth(); - int MAP_SZ4 = (OUT_WIDTH4 * problem.GetOutHeight() + read_unit - 1) / (read_unit); - // stride > 1 and/or apdding - if(problem.GetPadW() > 0 || problem.GetKernelStrideW() > 1 || problem.GetPadH() > 0 || - problem.GetKernelStrideH() > 1) - { - int step = is_forward ? read_unit : read_unit * problem.GetKernelStrideW(); - OUT_WIDTH4 = (problem.GetOutWidth() + step - 1) / (step); - int OUT_HEIGHT4 = is_forward - ? problem.GetOutHeight() - : (problem.GetOutHeight() + problem.GetKernelStrideH() - 1) / - problem.GetKernelStrideH(); - MAP_SZ4 = (OUT_WIDTH4 * OUT_HEIGHT4); - } - - int VERT_ALIGNED = 1; - int HORIZ_ALIGNED = 1; - if(!is_forward) - { - VERT_ALIGNED = - (problem.GetOutHeight() / problem.GetKernelStrideH() == problem.GetInHeight()) - ? 1 - : 0; - HORIZ_ALIGNED = - (problem.GetOutWidth() / problem.GetKernelStrideW() == problem.GetInWidth()) - ? 1 - : 0; - } - - int GRP_SZ = result.grp_tile0; - - // number of inputs inside wk-items - result.n_in_data_tiles = - std::min(static_cast(problem.GetInChannels()), result.n_in_data_tiles); - while(problem.GetInChannels() % result.n_in_data_tiles != 0 && - result.n_in_data_tiles > 1) - { - result.n_in_data_tiles /= 2; - } - - int CLOOP0 = - (problem.GetInChannels() + result.n_in_data_tiles - 1) / result.n_in_data_tiles; - - // number of outputs inside wk_item - result.n_out_pix_tiles = - std::min(static_cast(problem.GetOutChannels()), result.n_out_pix_tiles); - while(problem.GetOutChannels() % result.n_out_pix_tiles != 0 && - result.n_out_pix_tiles > 1) - { - result.n_out_pix_tiles /= 2; - } - - KernelInfo kernel; - - kernel.comp_options = - std::string(" -DMLO_DIR_FORWARD=") + (is_forward ? "1" : "0") + - std::string(" -DMLO_FILTER_SIZE0=") + std::to_string(problem.GetWeightsWidth()) + - std::string(" -DMLO_FILTER_SIZE1=") + std::to_string(problem.GetWeightsHeight()) + - std::string(" -DMLO_FILTER_STRIDE0=") + std::to_string(problem.GetKernelStrideW()) + - std::string(" -DMLO_FILTER_STRIDE1=") + std::to_string(problem.GetKernelStrideH()) + - std::string(" -DMLO_FILTER_PAD0=") + std::to_string(problem.GetPadW()) + - std::string(" -DMLO_FILTER_PAD1=") + std::to_string(problem.GetPadH()) + - std::string(" -DMLO_IN_WIDTH=") + std::to_string(problem.GetInWidth()) + - std::string(" -DMLO_IN_HEIGHT=") + std::to_string(problem.GetInHeight()) + - std::string(" -DMLO_OUT_WIDTH=") + std::to_string(problem.GetOutWidth()) + - std::string(" -DMLO_OUT_HEIGHT=") + std::to_string(problem.GetOutHeight()) + - std::string(" -DMLO_N_OUTPUTS=") + std::to_string(problem.GetOutChannels()) + - std::string(" -DMLO_N_INPUTS=") + std::to_string(problem.GetInChannels()) + - std::string(" -DMLO_BATCH_SZ=") + std::to_string(problem.GetBatchSize()) + - std::string(" -DMLO_OUT_BATCH_STRIDE=") + - std::to_string(problem.GetOutBatchStride()) + - std::string(" -DMLO_OUT_CHANNEL_STRIDE=") + - std::to_string(problem.GetOutChannelStride()) + std::string(" -DMLO_OUT_STRIDE=") + - std::to_string(problem.GetOutStrideH()) + std::string(" -DMLO_IN_BATCH_STRIDE=") + - std::to_string(problem.GetInBatchStride()) + - std::string(" -DMLO_IN_CHANNEL_STRIDE=") + - std::to_string(problem.GetInChannelStride()) + std::string(" -DMLO_IN_STRIDE=") + - std::to_string(problem.GetInStrideH()) + std::string(" -DMLO_WEI_BSTRIDE=") + - std::to_string(wei_bstride) + std::string(" -DMLO_WEI_CHANNEL_STRIDE=") + - std::to_string(wei_cstride) + - // algorithm parameters - std::string(" -DMLO_GRP_SZ0=") + std::to_string(GRP_SZ) + - std::string(" -DMLO_GRP_SZ1=") + std::to_string(1) + - std::string(" -DMLO_GRP_SZ2=") + std::to_string(1) + - - std::string(" -DMLO_MAP_SZ4=") + std::to_string(MAP_SZ4) + - std::string(" -DMLO_OUT_WIDTH4=") + std::to_string(OUT_WIDTH4) + - std::string(" -DMLO_VERT_ALIGNED=") + std::to_string(VERT_ALIGNED) + - std::string(" -DMLO_HORIZ_ALIGNED=") + std::to_string(HORIZ_ALIGNED) + - - std::string(" -DMLO_N_LCL_BATCHS=") + - std::to_string(result.n_stacks) + // # of diff stacks (part of batch). - std::string(" -DMLO_N_LCL_OUT_MAPS=") + - std::to_string(result.n_out_pix_tiles) + // # output pixel tiles per wk-item (ALU) - std::string(" -DMLO_N_LCL_IN_MAPS=") + - std::to_string( - result.n_in_data_tiles) + // total # of blocks of different inputs in LDS - std::string(" -DMLO_CONV_BIAS=") + - std::to_string(problem.GetBias()) + - - std::string(" -DMLO_READ_UNIT=") + std::to_string(read_unit) + - std::string(" -DMLO_CLOOP0=") + std::to_string(CLOOP0) + - - ctx.general_compile_options; - - kernel.l_wk.push_back(result.grp_tile0); - kernel.l_wk.push_back(1); - kernel.l_wk.push_back(1); - - size_t gbl_wk0 = problem.GetBatchSize() * MAP_SZ4; - - size_t gbl_wk1 = - (problem.GetOutChannels() + result.n_out_pix_tiles - 1) / result.n_out_pix_tiles; - size_t gbl_wk2 = 1; - - kernel.g_wk.push_back(gbl_wk0); - kernel.g_wk.push_back(gbl_wk1); - kernel.g_wk.push_back(gbl_wk2); - - kernel.kernel_file = "MIOpenConv1x1S.cl"; - kernel.kernel_name = - (problem.GetKernelStrideW() == 1 && problem.GetKernelStrideH() == 1) - ? "MIOpenConv1x1" - : "MIOpenConv1x1pquv"; - result.construction_params.push_back(kernel); - } - } - - result.invoker_factory = &miopen::conv::MakeGenericXWYPadInvoker; - return result; -} - -} // namespace conv -} // namespace solver -} // namespace miopen diff --git a/projects/miopen/src/solver/conv/conv_ocl_dir2Dfwd_exhaustive_search.cpp b/projects/miopen/src/solver/conv/conv_ocl_dir2Dfwd_exhaustive_search.cpp index 28665783a8fa..08190a0c2658 100644 --- a/projects/miopen/src/solver/conv/conv_ocl_dir2Dfwd_exhaustive_search.cpp +++ b/projects/miopen/src/solver/conv/conv_ocl_dir2Dfwd_exhaustive_search.cpp @@ -426,16 +426,15 @@ ConvOclDirectFwdLegacyExhaustiveSearch::SearchImpl(const ExecutionContext& ctx, result.n_in_data_tiles = (1 << i_t); } - const auto ret = - MeasurePerfConfig(profile_h, - bot_ocl_ptr, - top_ocl_ptr, - wei_ocl_ptr, - bias_ocl_ptr, - processing_time, - ctx, - problem, - result); + const auto ret = MeasurePerfConfig(profile_h, + bot_ocl_ptr, + top_ocl_ptr, + wei_ocl_ptr, + bias_ocl_ptr, + processing_time, + ctx, + problem, + result); --runs_left; if(ret != 0) { @@ -647,32 +646,15 @@ ConvOclDirectFwdLegacyExhaustiveSearch::SearchImpl(const ExecutionContext& ctx, int ret = -1; double default_time = std::numeric_limits::max(); const auto default_config = GetDefaultPerformanceConfig(ctx, problem); - if(problem.GetWeightsWidth() == 1 && problem.GetWeightsHeight() == 1 && - problem.GetGroupCount() == - 1) // Group conv: None 1x1 version yet, fallback to universal kernel. - { - ret = MeasurePerfConfig(profile_h, - bot_ocl_ptr, - top_ocl_ptr, - wei_ocl_ptr, - bias_ocl_ptr, - default_time, - ctx, - problem, - default_config); - } - else - { - ret = MeasurePerfConfig(profile_h, - bot_ocl_ptr, - top_ocl_ptr, - wei_ocl_ptr, - bias_ocl_ptr, - default_time, - ctx, - problem, - default_config); - } + ret = MeasurePerfConfig(profile_h, + bot_ocl_ptr, + top_ocl_ptr, + wei_ocl_ptr, + bias_ocl_ptr, + default_time, + ctx, + problem, + default_config); if(ret == 0) { is_passed = true; diff --git a/projects/miopen/test/gtest/unit_FinInterface.cpp b/projects/miopen/test/gtest/unit_FinInterface.cpp index ac90626a2cdd..075afe36ef97 100644 --- a/projects/miopen/test/gtest/unit_FinInterface.cpp +++ b/projects/miopen/test/gtest/unit_FinInterface.cpp @@ -1,28 +1,5 @@ -/******************************************************************************* - * - * MIT License - * - * Copyright (c) 2024 Advanced Micro Devices, Inc. - * - * 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. - * - *******************************************************************************/ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT #include #include @@ -216,7 +193,6 @@ const auto& GetSolversInfo() {"ConvHipDirectFwd11x11", {8, false, false, "miopenConvolutionFwdAlgoDirect"}}, {"ConvOclDirectFwdGen", {9, false, false, "miopenConvolutionFwdAlgoDirect"}}, {"ConvOclDirectFwd", {11, false, true, "miopenConvolutionFwdAlgoDirect"}}, - {"ConvOclDirectFwd1x1", {13, false, true, "miopenConvolutionFwdAlgoDirect"}}, {"ConvBinWinograd3x3U", {14, true, false, "miopenConvolutionFwdAlgoWinograd"}}, {"ConvBinWinogradRxS", {15, true, false, "miopenConvolutionFwdAlgoWinograd"}}, {"ConvAsmBwdWrW3x3", {16, false, true, "miopenConvolutionFwdAlgoDirect"}}, diff --git a/projects/miopen/test/gtest/unit_conv_solver_ConvOclDirectFwd1x1.cpp b/projects/miopen/test/gtest/unit_conv_solver_ConvOclDirectFwd1x1.cpp deleted file mode 100644 index 7ddc715a25c4..000000000000 --- a/projects/miopen/test/gtest/unit_conv_solver_ConvOclDirectFwd1x1.cpp +++ /dev/null @@ -1,217 +0,0 @@ -/******************************************************************************* - * - * MIT License - * - * Copyright (c) 2024 Advanced Micro Devices, Inc. - * - * 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. - * - *******************************************************************************/ - -#include "unit_conv_solver.hpp" - -MIOPEN_LIB_ENV_VAR(MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1) - -namespace { - -class WA_SWDEV_271887_ScopedDisabler -{ -public: - WA_SWDEV_271887_ScopedDisabler() - { - if(MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1) - prev = lib_env::value(MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1); - if(prev != true) - lib_env::update(MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1, true); - } - - ~WA_SWDEV_271887_ScopedDisabler() - { - if(prev) - { - if(prev != true) - lib_env::update(MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1, false); - } - else - { - lib_env::clear(MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1); - } - } - -private: - std::optional prev; -}; - -auto GetConvTestCases(miopenDataType_t datatype) -{ - using TestCase = miopen::unit_tests::ConvTestCase; - - return std::vector{ - // clang-format off - TestCase{{16, 16, 16, 16}, {16, 16, 1, 1}, {0, 0}, {1, 1}, {1, 1}, datatype}, - // clang-format on - }; -} - -auto GetConvTestCasesFull(miopenDataType_t datatype, miopen::conv::Direction direction) -{ - using TestCase = miopen::unit_tests::ConvTestCase; - - auto cases = std::vector{}; - - if(datatype == miopenHalf && direction == miopen::conv::Direction::Forward) - { - // clang-format off - // Regression test for https://github.com/ROCm/MIOpen/issues/894 - cases.emplace_back(TestCase{{1, 16, 7, 7}, {16, 16, 1, 1}, {0, 0}, {1, 1}, {1, 1}, miopenHalf}); - // clang-format on - } - - return cases; -} - -const auto& GetTestParams() -{ - static const auto params = [] { - Gpu supported_gpus = Gpu::gfx900 | Gpu::gfx906 | Gpu::gfx908 | Gpu::gfx90A | Gpu::gfx103X; - auto p = miopen::unit_tests::UnitTestConvSolverParams(supported_gpus); - p.Tunable(5); - return p; - }(); - return params; -} - -const auto& GetTestParamsFull() -{ - static const auto params = [] { - Gpu supported_gpus = Gpu::gfx900 | Gpu::gfx906 | Gpu::gfx908 | Gpu::gfx90A | Gpu::gfx103X; - auto p = miopen::unit_tests::UnitTestConvSolverParams(supported_gpus); - p.Tunable(1000); - return p; - }(); - return params; -} - -} // namespace - -using GPU_UnitTestConvSolverOclDirectFwd1x1Fwd_FP16 = GPU_UnitTestConvSolverFwd_FP16; -using GPU_UnitTestConvSolverOclDirectFwd1x1Bwd_FP16 = GPU_UnitTestConvSolverBwd_FP16; - -using GPU_UnitTestConvSolverOclDirectFwd1x1Fwd_BFP16 = GPU_UnitTestConvSolverFwd_BFP16; -using GPU_UnitTestConvSolverOclDirectFwd1x1Bwd_BFP16 = GPU_UnitTestConvSolverBwd_BFP16; - -using GPU_UnitTestConvSolverOclDirectFwd1x1Fwd_FP32 = GPU_UnitTestConvSolverFwd_FP32; -using GPU_UnitTestConvSolverOclDirectFwd1x1Bwd_FP32 = GPU_UnitTestConvSolverBwd_FP32; - -using CPU_UnitTestConvSolverOclDirectFwd1x1DevApplicabilityFwd_NONE = - CPU_UnitTestConvSolverDevApplicabilityFwd_NONE; - -TEST_P(GPU_UnitTestConvSolverOclDirectFwd1x1Fwd_FP16, ConvOclDirectFwd1x1) -{ - WA_SWDEV_271887_ScopedDisabler wa_swdev_271887_disabler; - this->RunTest(miopen::solver::conv::ConvOclDirectFwd1x1{}); -}; - -TEST_P(GPU_UnitTestConvSolverOclDirectFwd1x1Bwd_FP16, ConvOclDirectFwd1x1) -{ - WA_SWDEV_271887_ScopedDisabler wa_swdev_271887_disabler; - this->RunTest(miopen::solver::conv::ConvOclDirectFwd1x1{}); -}; - -TEST_P(GPU_UnitTestConvSolverOclDirectFwd1x1Fwd_BFP16, ConvOclDirectFwd1x1) -{ - WA_SWDEV_271887_ScopedDisabler wa_swdev_271887_disabler; - this->RunTest(miopen::solver::conv::ConvOclDirectFwd1x1{}); -}; - -TEST_P(GPU_UnitTestConvSolverOclDirectFwd1x1Bwd_BFP16, ConvOclDirectFwd1x1) -{ - WA_SWDEV_271887_ScopedDisabler wa_swdev_271887_disabler; - this->RunTest(miopen::solver::conv::ConvOclDirectFwd1x1{}); -}; - -TEST_P(GPU_UnitTestConvSolverOclDirectFwd1x1Fwd_FP32, ConvOclDirectFwd1x1) -{ - WA_SWDEV_271887_ScopedDisabler wa_swdev_271887_disabler; - this->RunTest(miopen::solver::conv::ConvOclDirectFwd1x1{}); -}; - -TEST_P(GPU_UnitTestConvSolverOclDirectFwd1x1Bwd_FP32, ConvOclDirectFwd1x1) -{ - WA_SWDEV_271887_ScopedDisabler wa_swdev_271887_disabler; - this->RunTest(miopen::solver::conv::ConvOclDirectFwd1x1{}); -}; - -TEST_P(CPU_UnitTestConvSolverOclDirectFwd1x1DevApplicabilityFwd_NONE, ConvOclDirectFwd1x1) -{ - WA_SWDEV_271887_ScopedDisabler wa_swdev_271887_disabler; - this->RunTest(miopen::solver::conv::ConvOclDirectFwd1x1{}); -}; - -// Smoke tests -INSTANTIATE_TEST_SUITE_P(Smoke, - GPU_UnitTestConvSolverOclDirectFwd1x1Fwd_FP16, - testing::Combine(testing::Values(GetTestParams()), - testing::Values(miopenConvolutionAlgoDirect), - testing::ValuesIn(GetConvTestCases(miopenHalf)))); - -INSTANTIATE_TEST_SUITE_P(Smoke, - GPU_UnitTestConvSolverOclDirectFwd1x1Bwd_FP16, - testing::Combine(testing::Values(GetTestParams()), - testing::Values(miopenConvolutionAlgoDirect), - testing::ValuesIn(GetConvTestCases(miopenHalf)))); - -INSTANTIATE_TEST_SUITE_P(Smoke, - GPU_UnitTestConvSolverOclDirectFwd1x1Fwd_BFP16, - testing::Combine(testing::Values(GetTestParams()), - testing::Values(miopenConvolutionAlgoDirect), - testing::ValuesIn(GetConvTestCases(miopenBFloat16)))); - -INSTANTIATE_TEST_SUITE_P(Smoke, - GPU_UnitTestConvSolverOclDirectFwd1x1Bwd_BFP16, - testing::Combine(testing::Values(GetTestParams()), - testing::Values(miopenConvolutionAlgoDirect), - testing::ValuesIn(GetConvTestCases(miopenBFloat16)))); - -INSTANTIATE_TEST_SUITE_P(Smoke, - GPU_UnitTestConvSolverOclDirectFwd1x1Fwd_FP32, - testing::Combine(testing::Values(GetTestParams()), - testing::Values(miopenConvolutionAlgoDirect), - testing::ValuesIn(GetConvTestCases(miopenFloat)))); - -INSTANTIATE_TEST_SUITE_P(Smoke, - GPU_UnitTestConvSolverOclDirectFwd1x1Bwd_FP32, - testing::Combine(testing::Values(GetTestParams()), - testing::Values(miopenConvolutionAlgoDirect), - testing::ValuesIn(GetConvTestCases(miopenFloat)))); - -// Device applicability test -INSTANTIATE_TEST_SUITE_P(Smoke, - CPU_UnitTestConvSolverOclDirectFwd1x1DevApplicabilityFwd_NONE, - testing::Combine(testing::Values(GetTestParams()), - testing::Values(GetConvTestCases(miopenFloat)[0]))); - -// Full tests -// clang-format off -INSTANTIATE_TEST_SUITE_P(Full, - GPU_UnitTestConvSolverOclDirectFwd1x1Fwd_FP16, - testing::Combine(testing::Values(GetTestParamsFull()), - testing::Values(miopenConvolutionAlgoDirect), - testing::ValuesIn(GetConvTestCasesFull(miopenHalf, miopen::conv::Direction::Forward)))); -// clang-format on