diff --git a/CMakeLists.txt b/CMakeLists.txt index ac8484172b..aebd984989 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -68,6 +68,9 @@ include(ROCMCreatePackage) include(CheckCXXCompilerFlag) include(ROCMHeaderWrapper) +# Build library with Beta APIs +add_definitions("-DMIOPEN_BETA_API=1") + set(MIOPEN_ENABLE_AI_IMMED_MODE_FALLBACK On CACHE BOOL "Enable AI-based fallback for Immediate Mode") set(MIOPEN_ENABLE_AI_KERNEL_TUNING On CACHE BOOL "Enable AI heuristic for kernel tuning") set(MIOPEN_ENABLE_SQLITE On CACHE BOOL "") diff --git a/driver/layernorm_driver.hpp b/driver/layernorm_driver.hpp index 8251472625..e43f3f2f37 100644 --- a/driver/layernorm_driver.hpp +++ b/driver/layernorm_driver.hpp @@ -255,19 +255,19 @@ int LayerNormDriver::AllocateBuffersAndCopy() for(int i = 0; i < in_sz; i++) { - in[i] = RAN_GEN(static_cast(0.0), static_cast(1.0)); + in[i] = prng::gen_A_to_B(static_cast(0.0), static_cast(1.0)); } status = in_dev->ToGPU(q, in.data()); for(int i = 0; i < weight_sz; i++) { - weight[i] = RAN_GEN(static_cast(0.0), static_cast(1.0)); + weight[i] = prng::gen_A_to_B(static_cast(0.0), static_cast(1.0)); } status = weight_dev->ToGPU(q, weight.data()); for(int i = 0; i < bias_sz; i++) { - bias[i] = RAN_GEN(static_cast(0.0), static_cast(1.0)); + bias[i] = prng::gen_A_to_B(static_cast(0.0), static_cast(1.0)); } status = bias_dev->ToGPU(q, bias.data()); diff --git a/fin b/fin index b2f3f4db3c..26b5c32864 160000 --- a/fin +++ b/fin @@ -1 +1 @@ -Subproject commit b2f3f4db3c3d7dd757e6d9e68719a780d8114dfa +Subproject commit 26b5c328642a6af5041539ceae36b9340829384b diff --git a/include/miopen/miopen.h b/include/miopen/miopen.h index a59181acf3..abbec2599c 100644 --- a/include/miopen/miopen.h +++ b/include/miopen/miopen.h @@ -112,11 +112,13 @@ typedef enum miopenStatusVersionMismatch = 10, /*!< Version mismatch of the supplied binary data argment. */ } miopenStatus_t; +#ifdef MIOPEN_BETA_API typedef enum { miopenF8RoundingModeStandard = 0, miopenF8RoundingModeStochastic = 1, } miopenF8RoundingMode_t; +#endif /*! @brief Get character string for an error code. * @@ -354,9 +356,14 @@ typedef enum 4, /*!< Pack of four 8-bit int points in NCHW_VECT_C format (Partially supported) */ miopenBFloat16 = 5, /*!< 16-bit binary floating point (8-bit exponent, 7-bit fraction) (Partially supported) */ - miopenDouble = 6, /*!< 64-bit floating point (Partially supported) */ + miopenDouble = 6, /*!< 64-bit floating point (Partially supported) */ +#ifdef MIOPEN_BETA_API miopenFloat8 = 7, - miopenBFloat8 = 8 + miopenBFloat8 = 8, +#else +// miopenReserved1 = 7, +// miopenReserved2 = 8, +#endif } miopenDataType_t; /*! @ingroup tensor @@ -601,11 +608,15 @@ typedef enum MIOPEN_CONVOLUTION_ATTRIB_DETERMINISTIC = 1, /*!< Restrict MIOpen convolutions to kernels which produce numerically deterministic results. 0 - disabled (default), 1 - enabled >*/ +#ifdef MIOPEN_BETA_API MIOPEN_CONVOLUTION_ATTRIB_FP8_ROUNDING_MODE = 2, /*!*/ +#else +// miopenReserved1 = 2, +#endif } miopenConvolutionAttrib_t; /** @addtogroup tensor @@ -723,6 +734,7 @@ MIOPEN_EXPORT miopenStatus_t miopenSetTensorDescriptor(miopenTensorDescriptor_t const int* dimsA, const int* stridesA); +#ifdef MIOPEN_BETA_API /*! @brief Set the tensor cast type * * For tensors where the cast_type attribute is set, the tensor elements would be converted to the @@ -734,6 +746,7 @@ MIOPEN_EXPORT miopenStatus_t miopenSetTensorDescriptor(miopenTensorDescriptor_t */ MIOPEN_EXPORT miopenStatus_t miopenSetTensorCastType(miopenTensorDescriptor_t tensorDesc, miopenDataType_t cast_type); +#endif /*! @brief Set shape of N-dimensional tensor * diff --git a/src/include/miopen/layernorm.hpp b/src/include/miopen/layernorm.hpp index 8ec2d96055..f897e79eea 100644 --- a/src/include/miopen/layernorm.hpp +++ b/src/include/miopen/layernorm.hpp @@ -49,8 +49,8 @@ miopenStatus_t LayerNormForward(const Handle& handle, const TensorDescriptor& rstdDesc, Data_t rstd, miopenLayerNormMode_t mode, - const float epsilon, - const int32_t normalized_dim); + float epsilon, + int32_t normalized_dim); } // namespace miopen #endif // _MIOPEN_LAYERNORM_HPP_ diff --git a/src/layer_norm.cpp b/src/layer_norm.cpp index 3d52bc771f..33030887ee 100644 --- a/src/layer_norm.cpp +++ b/src/layer_norm.cpp @@ -48,8 +48,8 @@ miopenStatus_t LayerNormForward(const Handle& handle, const TensorDescriptor& rstdDesc, Data_t rstd, miopenLayerNormMode_t mode, - const float epsilon, - const int32_t normalized_dim) + float epsilon, + int32_t normalized_dim) { if(x == nullptr || y == nullptr) { diff --git a/src/solver/batchnorm/backward_ck.cpp b/src/solver/batchnorm/backward_ck.cpp index fba8724990..6fecc24049 100644 --- a/src/solver/batchnorm/backward_ck.cpp +++ b/src/solver/batchnorm/backward_ck.cpp @@ -156,11 +156,11 @@ static bool CheckCKApplicability(const miopen::batchnorm::ProblemDescription& pr #endif -bool BnCKBwdBackward::IsApplicable(const ExecutionContext& ctx, +bool BnCKBwdBackward::IsApplicable(const ExecutionContext& context, const miopen::batchnorm::ProblemDescription& bn_problem) const { #if !MIOPEN_BACKEND_HIP || !MIOPEN_USE_COMPOSABLEKERNEL - std::ignore = ctx; + std::ignore = context; std::ignore = fdesc_problem; return false; #else @@ -168,7 +168,7 @@ bool BnCKBwdBackward::IsApplicable(const ExecutionContext& ctx, return false; if(!bn_problem.IsLayoutNHWC()) return false; - if(!ck_utility::is_ck_supported_hardware(ctx.GetStream())) + if(!ck_utility::is_ck_supported_hardware(context.GetStream())) return false; if(bn_problem.GetXDesc().GetType() != bn_problem.GetScaleBiasDiffDesc().GetType()) return false;