[MIOpen] -Wno-conversion

projects/miopen/addkernels/addkernels.cpp

@@ -41,16 +41,16 @@ void Bin2Hex(std::istream& source,
 
     while(blockStart < sourceSize)
     {
-        source.read(reinterpret_cast<char*>(buffer.get()), bufferSize);
+        source.read(reinterpret_cast<char*>(buffer.get()), std::streamsize(bufferSize));
 
-        const std::streamoff pos       = source.tellg();
-        const std::streamoff blockSize = (pos < 0 ? sourceSize : pos) - blockStart;
-        std::streamoff i               = 0;
+        const std::streamoff pos = source.tellg();
+        const auto blockSize = (pos < 0 ? sourceSize : pos) - blockStart;
 
+        size_t i = 0;
         while(i < blockSize)
         {
-            size_t j         = i;
-            const size_t end = std::min<size_t>(i + lineSize, blockSize);
+            size_t j = i;
+            const auto end = std::min<size_t>(i + lineSize, size_t(blockSize));
 
             for(; j < end; j++)
                 target << "0x" << std::setw(2) << static_cast<unsigned>(buffer[j]) << ",";
@@ -98,7 +98,7 @@ void Bin2Asm(std::istream& source,
     // Write binary data
     for(std::streamoff blockStart = 0; blockStart < sourceSize; blockStart += bufferSize)
     {
-        source.read(buffer.get(), bufferSize);
+        source.read(buffer.get(), std::streamsize(bufferSize));
 
         const auto pos       = source.tellg();
         const auto blockSize = (pos < 0 ? sourceSize : pos) - blockStart;
@@ -329,11 +329,11 @@ int main(int argc, char* argv[])
         }
         else if(arg == "-l" || arg == "-line-size")
         {
-            lineSize = std::stol(argv[++i]);
+            lineSize = std::stoul(argv[++i]);
         }
         else if(arg == "-b" || arg == "-buffer")
         {
-            bufferSize = std::stol(argv[++i]);
+            bufferSize = std::stoul(argv[++i]);
         }
         else if(arg == "-g" || arg == "-guard")
         {
@@ -446,7 +446,7 @@ int main(int argc, char* argv[])
             std::cerr << "failure opening file: " << targetFile << "\n";
             return 1;
         }
-        file.write(sourceCode.data(), sourceCode.length());
+        file.write(sourceCode.data(), std::streamsize(sourceCode.length()));
     }
 
     return 0;

projects/miopen/addkernels/include_inliner.cpp

@@ -89,7 +89,7 @@ void IncludeInliner::ProcessCore(std::istream& input,
 
         if(!word.empty() && word == directive && recurse)
         {
-            auto first_quote_pos = line.find('"', static_cast<int>(line_parser.tellg()) + 1);
+            auto first_quote_pos = line.find('"', static_cast<size_t>(line_parser.tellg()) + 1);
             std::string::size_type second_quote_pos;
 
             if(first_quote_pos != std::string::npos)
@@ -103,7 +103,7 @@ void IncludeInliner::ProcessCore(std::istream& input,
                 if(!allow_angle_brackets)
                     throw WrongInlineDirectiveException(GetIncludeStackTrace(current_line));
 
-                first_quote_pos = line.find('<', static_cast<int>(line_parser.tellg()) + 1);
+                first_quote_pos = line.find('<', static_cast<size_t>(line_parser.tellg()) + 1);
                 if(first_quote_pos == std::string::npos)
                     throw WrongInlineDirectiveException(GetIncludeStackTrace(current_line));
 

projects/miopen/cmake/EnableCompilerWarnings.cmake

@@ -46,7 +46,6 @@ set(__clang_cxx_compile_options
     -Weverything
     -Wno-c++98-compat
     -Wno-c++98-compat-pedantic
-    -Wno-conversion
     -Wno-double-promotion
     -Wno-exit-time-destructors
     -Wno-extra-semi
@@ -71,7 +70,6 @@ set(__clang_cxx_compile_options
 
 if(CMAKE_CXX_COMPILER_ID STREQUAL "Clang" AND CMAKE_CXX_COMPILER_VERSION VERSION_GREATER_EQUAL "19")
     list(APPEND __clang_cxx_compile_options
-        -Wno-unique-object-duplication
         -Wno-switch-default)
 endif()
 
@@ -85,6 +83,8 @@ endif()
 
 if(WIN32)
     list(APPEND __clang_cxx_compile_options
+        -Wno-ignored-attributes
+        -Wno-language-extension-token
         -fms-extensions
         -fms-compatibility)
 endif()

projects/miopen/driver/layernorm_driver.hpp

@@ -498,8 +498,9 @@ int LayerNormDriver<T>::RunBackwardGPU()
             std::cout << "Wall-clock Time Backward LayerNorm Elapsed: " << t.gettime_ms() / iter
                       << " ms\n";
 
-        float kernel_average_time =
-            iter > 1 ? (kernel_total_time - kernel_first_time) / (iter - 1) : kernel_first_time;
+        float kernel_average_time = iter > 1
+                                        ? (kernel_total_time - kernel_first_time) / float(iter - 1)
+                                        : kernel_first_time;
         std::cout << "GPU Kernel Time Backward LayerNorm Elapsed: " << kernel_average_time
                   << " ms\n";
     }

projects/miopen/driver/timer.hpp

@@ -47,7 +47,7 @@
 class Timer
 {
 public:
-    Timer(){};
+    Timer() {};
     void start(const bool enabled = true)
     {
         if(!enabled)
@@ -80,7 +80,7 @@ class Timer
 class Timer2
 {
 public:
-    Timer2(){};
+    Timer2() {};
     void start(const bool enabled = true)
     {
         if(!enabled)
@@ -262,7 +262,7 @@ class RNNCombTimeLogger
 
         printf("GPU Kernel Time Elapsed: %f ms\n", n_iter > 1 ? gpu_avg / (n_iter - 1) : gpu_time);
         printf("Wall-clock Time Elapsed: %f ms\n",
-               n_iter > 1 ? host_avg / (n_iter - 1) : hostTimePerLaunch[0]);
+               n_iter > 1 ? host_avg / float(n_iter - 1) : hostTimePerLaunch[0]);
     }
 
     enum class ClockMode

projects/miopen/src/activ/problem_description.cpp

@@ -54,7 +54,7 @@ NetworkConfig ProblemDescription::MakeNetworkConfig() const
 
     const auto read_len = (packed) ? x_elem_sz : x_width2D;
 
-    const auto read_unit = (read_len % 4 == 0) ? 4 : (read_len % 2 == 0) ? 2 : 1;
+    const auto read_unit = (read_len % 4u == 0) ? 4u : (read_len % 2u == 0) ? 2u : 1u;
     const auto MAP_RD    = read_len / read_unit;
 
     std::ostringstream ss;

projects/miopen/src/api/find2_0_commons.cpp

@@ -272,7 +272,7 @@ miopenStatus_t miopenFindSolutions(miopenHandle_t handle,
             },
             problem_deref);
 
-        for(auto i = 0; i < solutions_deref.size(); ++i)
+        for(size_t i = 0; i < solutions_deref.size(); ++i)
         {
             auto& theSolution = miopen::deref(solutions + i);
             theSolution       = new miopen::Solution{std::move(solutions_deref[i])};

projects/miopen/src/base64.cpp

@@ -55,8 +55,8 @@ std::string base64Encode(const uint8_t* data, std::size_t length)
         if(i == 3)
         {
             charArray4[0] = (charArray3[0] & 0xfc) >> 2;
-            charArray4[1] = ((charArray3[0] & 0x03) << 4) + ((charArray3[1] & 0xf0) >> 4);
-            charArray4[2] = ((charArray3[1] & 0x0f) << 2) + ((charArray3[2] & 0xc0) >> 6);
+            charArray4[1] = uint8_t((charArray3[0] & 0x03) << 4) + uint8_t((charArray3[1] & 0xf0) >> 4);
+            charArray4[2] = uint8_t((charArray3[1] & 0x0f) << 2) + uint8_t((charArray3[2] & 0xc0) >> 6);
             charArray4[3] = charArray3[2] & 0x3f;
 
             std::transform(std::begin(charArray4),
@@ -75,8 +75,8 @@ std::string base64Encode(const uint8_t* data, std::size_t length)
         }
 
         charArray4[0] = (charArray3[0] & 0xfc) >> 2;
-        charArray4[1] = ((charArray3[0] & 0x03) << 4) + ((charArray3[1] & 0xf0) >> 4);
-        charArray4[2] = ((charArray3[1] & 0x0f) << 2) + ((charArray3[2] & 0xc0) >> 6);
+        charArray4[1] = uint8_t((charArray3[0] & 0x03) << 4) + uint8_t((charArray3[1] & 0xf0) >> 4);
+        charArray4[2] = uint8_t((charArray3[1] & 0x0f) << 2) + uint8_t((charArray3[2] & 0xc0) >> 6);
         charArray4[3] = charArray3[2] & 0x3f;
 
         for(size_t j = 0; j < i + 1; j++)
@@ -123,9 +123,9 @@ std::vector<uint8_t> base64Decode(const std::string_view& encodedString)
         charArray4[i++] = static_cast<uint8_t>(std::distance(base64Chars.begin(), it));
         if(i == 4)
         {
-            charArray3[0] = (charArray4[0] << 2) + ((charArray4[1] & 0x30) >> 4);
-            charArray3[1] = ((charArray4[1] & 0xf) << 4) + ((charArray4[2] & 0x3c) >> 2);
-            charArray3[2] = ((charArray4[2] & 0x3) << 6) + charArray4[3];
+            charArray3[0] = uint8_t(charArray4[0] << 2) + uint8_t((charArray4[1] & 0x30) >> 4);
+            charArray3[1] = uint8_t((charArray4[1] & 0xf) << 4) + uint8_t((charArray4[2] & 0x3c) >> 2);
+            charArray3[2] = uint8_t((charArray4[2] & 0x3) << 6) + charArray4[3];
 
             decodedData.insert(decodedData.end(), charArray3, charArray3 + 3);
             i = 0;
@@ -139,9 +139,9 @@ std::vector<uint8_t> base64Decode(const std::string_view& encodedString)
             charArray4[j] = 0;
         }
 
-        charArray3[0] = (charArray4[0] << 2) + ((charArray4[1] & 0x30) >> 4);
-        charArray3[1] = ((charArray4[1] & 0xf) << 4) + ((charArray4[2] & 0x3c) >> 2);
-        charArray3[2] = ((charArray4[2] & 0x3) << 6) + charArray4[3];
+        charArray3[0] = uint8_t(charArray4[0] << 2) + uint8_t((charArray4[1] & 0x30) >> 4);
+        charArray3[1] = uint8_t((charArray4[1] & 0xf) << 4) + uint8_t((charArray4[2] & 0x3c) >> 2);
+        charArray3[2] = uint8_t((charArray4[2] & 0x3) << 6) + charArray4[3];
 
         for(size_t j = 0; j < i - 1; j++)
         {

projects/miopen/src/buffer_info.cpp

@@ -107,11 +107,16 @@ MemLayout_t GetGroupConvLayout(MemLayout_t layout, bool IsDataBuffer)
     MIOPEN_THROW(std::string("Internal error in GetGroupConvLayout: Unknown MemLayout_t "));
 }
 
-BuffInfo::BuffInfo(MemLayout_t layout, int nk, int c, int h, int w, int g, int _element_size)
+BuffInfo::BuffInfo(MemLayout_t layout,
+                   unsigned nk,
+                   unsigned c,
+                   unsigned h,
+                   unsigned w,
+                   unsigned g,
+                   size_t _element_size)
 {
-
     element_size       = _element_size;
-    const size_t count = static_cast<size_t>(nk) * c * h * w * g;
+    const size_t count = static_cast<size_t>(nk * c * h * w * g);
     total_byte_size    = count * element_size;
     size.nk            = nk;
     size.g             = g;

projects/miopen/src/cat/problem_description.cpp

@@ -46,13 +46,13 @@ NetworkConfig ProblemDescription::MakeNetworkConfig() const
     for(int i = 0; i < xCount; i++)
     {
         auto xlength   = xDescs[i]->GetLengths();
-        max_x_dim_size = std::max(max_x_dim_size, xlength[dim]);
+        max_x_dim_size = std::max(max_x_dim_size, xlength[size_t(dim)]);
     }
 
     auto ylength    = yDesc.GetLengths();
     auto outer_size = std::accumulate(
         ylength.begin(), ylength.begin() + dim, static_cast<size_t>(1), std::multiplies<size_t>());
-    auto stride         = yDesc.GetStrides()[dim];
+    auto stride         = yDesc.GetStrides()[size_t(dim)];
     auto dtype          = yDesc.GetType();
     auto data_size      = get_data_size(dtype);
     auto max_inner_size = max_x_dim_size * stride * data_size / sizeof(short4);

projects/miopen/src/check_numerics.cpp

@@ -37,7 +37,7 @@ namespace miopen {
 
 bool CheckNumericsEnabled(const int bitMask)
 {
-    return (env::value(MIOPEN_CHECK_NUMERICS) & bitMask) != 0;
+    return (env::value(MIOPEN_CHECK_NUMERICS) & static_cast<unsigned>(bitMask)) != 0;
 }
 
 // Must keep this structure synchronized with one in MIOpenCheckNumerics
@@ -73,7 +73,7 @@ std::string GetKernelName(miopenDataType_t data_type)
 bool checkNumericsImpl(
     const Handle& handle, int mode, const TensorDescriptor& dDesc, ConstData_t data, bool isInput)
 {
-    int numElements = dDesc.GetElementSize();
+    size_t numElements = dDesc.GetElementSize();
     CheckNumericsResult abnormal_h;
     auto abnormal_d =
         handle.Create(sizeof(CheckNumericsResult)); // TODO - someday avoid slow malloc/free here
@@ -123,8 +123,8 @@ bool checkNumericsImpl(
         {
             assert(numElements != 0);
             MIOPEN_LOG((isAbnormal ? miopen::LoggingLevel::Warning : miopen::LoggingLevel::Info),
-                       "Stats: mean=" << (abnormal_h.sum / numElements)
-                                      << " absmean=" << (abnormal_h.absSum / numElements)
+                       "Stats: mean=" << (abnormal_h.sum / float(numElements))
+                                      << " absmean=" << (abnormal_h.absSum / float(numElements))
                                       << " min=" << abnormal_h.min << " max=" << abnormal_h.max);
         }
     }
@@ -158,15 +158,15 @@ bool checkNumericsImpl(
 // Returns: 1 if abnormal value (inf or nan) detected in specified data, 0 otherwise
 bool checkNumericsInput(const Handle& handle, const TensorDescriptor& dDesc, ConstData_t data)
 {
-    return checkNumericsImpl(handle, env::value(MIOPEN_CHECK_NUMERICS), dDesc, data, true);
+    return checkNumericsImpl(handle, (int)env::value(MIOPEN_CHECK_NUMERICS), dDesc, data, true);
 }
 
 // Synchronizes to wait for kernel to finish, then checks data for output:
 // Returns: 1 if abnormal value (inf or nan) detected in specified data, 0 otherwise
 bool checkNumericsOutput(const Handle& handle, const TensorDescriptor& dDesc, ConstData_t data)
 {
     handle.Finish();
-    return checkNumericsImpl(handle, env::value(MIOPEN_CHECK_NUMERICS), dDesc, data, false);
+    return checkNumericsImpl(handle, (int)env::value(MIOPEN_CHECK_NUMERICS), dDesc, data, false);
 }
 
 } // namespace miopen

projects/miopen/src/conv/invokers/gcn_asm_wino.cpp

@@ -113,13 +113,13 @@ InvokerFactory MakeGcnAsmWinoV2InvokerFactory(const WinoShaderArgsV2& args,
             uint64_t a_offset = 0;
 
             // activation parameters
-            float alpha = 0.0f;
-            float beta  = 0.0f;
+            double alpha = 0.0;
+            double beta  = 0.0;
 
             if(fused && (args.activation_mode != WinoShaderActivationModeV2_t::IDENTITY))
             {
                 const auto& invoke_ctx = primitive_params.CastTo<fusion::FusionInvokeParams>();
-                const int idx          = do_bias ? 2 : 1;
+                const size_t idx       = do_bias ? 2 : 1;
                 const auto& activ_args =
                     dynamic_cast<fusion::ActivationOpInvokeParam&>(*invoke_ctx.op_args.params[idx]);
                 if(args.activation_mode == WinoShaderActivationModeV2_t::SCALED_TANH)
@@ -187,8 +187,8 @@ InvokerFactory MakeGcnAsmWinoV2InvokerFactory(const WinoShaderArgsV2& args,
                 args.out_h,               // uint32_t,    output height
                 args.out_w,               // uint32_t,    output width
                 bias_addr,                // uint64_t,    address of bias buffer
-                alpha,                    // fp32,        activation parameter alpha
-                beta,                     // fp32,        activation parameter beta
+                float(alpha),             // fp32,        activation parameter alpha
+                float(beta),              // fp32,        activation parameter beta
                 d_offset,                 // uint64_t,    byte offset for buffer referenced by data_addr
                 f_offset,                 // uint64_t,    byte offset for buffer referenced by filter_addr
                 o_offset,                 // uint64_t,    byte offset for buffer referenced by output_addr

projects/miopen/src/conv/invokers/impl_gemm.cpp

@@ -53,10 +53,10 @@ InvokerFactory MakeImplGemmDataInvokerFactory(const ProblemDescription& problem)
                     bool need_atomic_add        = false;
                     bool every_pixel_is_written = true;
 
-                    for(int i = 0; i < conv.GetSpatialDimension(); ++i)
+                    for(size_t i = 0; i < conv.GetSpatialDimension(); ++i)
                     {
                         const auto conv_stride   = conv.GetConvStrides()[i];
-                        const auto conv_dilation = conv.GetConvDilations()[i];
+                        const auto conv_dilation = size_t(conv.GetConvDilations()[i]);
                         const auto filter_size   = tensors.wDesc.GetLengths()[2 + i];
 
                         if(conv_stride < conv_dilation * (filter_size - 1) + 1)
@@ -170,7 +170,7 @@ InvokerFactory MakeImplGemmDataInvokerFactory(const ProblemDescription& problem)
                 {
                     bool every_pixel_is_written = true;
 
-                    for(int i = 0; i < conv.GetSpatialDimension(); ++i)
+                    for(size_t i = 0; i < conv.GetSpatialDimension(); ++i)
                     {
                         const auto conv_stride   = conv.GetConvStrides()[i];
                         const auto conv_dilation = conv.GetConvDilations()[i];