Clean up warnings (NVIDIA#5811)

Signed-off-by: Joaquin Anton Guirao <[email protected]>

Author: jantonguirao

CMakeLists.txt

@@ -286,6 +286,15 @@ message(STATUS "Generated CMAKE_CUDA_ARCHITECTURES: ${CMAKE_CUDA_ARCHITECTURES}"
 
 set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --compiler-options \"-fvisibility=hidden -Wno-free-nonheap-object\" --Wno-deprecated-gpu-targets -Xfatbin -compress-all")
 
+# Suppress a few warnings from nvcc:
+# calling a __host__ function(...) from a __host__ __device__ function(...) is not allowed
+set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --diag-suppress=20011")
+# 'long double' is treated as 'double' in device code
+set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --diag-suppress=20208")
+# unrecognized #pragma in device code
+set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --diag-suppress=20199")
+# overloaded virtual function "dali::Operator<dali::CPUBackend>::RunImpl" is only partially overridden in class "..."
+set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --diag-suppress=611")
 
 # Dependencies
 if(${ARCH} MATCHES "aarch64-linux")

dali/kernels/slice/slice_hwc2chw_normalize_gpu.cu

@@ -130,7 +130,7 @@ __device__ __forceinline__ Tile *load_linear_tile(Tile *tile,
     aligned_tile[idx * 4 + 3] = in.w;
   }
 
-  uint32_t processed_in_main = left_after_prologue & -4;  // equivalent to (x / 4) * 4
+  uint32_t processed_in_main = left_after_prologue & ~0x3;;  // equivalent to (x / 4) * 4
   uint32_t left_after_main = left_after_prologue - processed_in_main;
 
   // epilogue
@@ -238,7 +238,7 @@ __device__ __forceinline__ Tile *slice_load_linear_tile(
       aligned_tile[idx * 4 + 3] = in.w;
     }
 
-    uint32_t processed_in_main = left_after_prologue & -4;  // equivalent to (x / 4) * 4
+    uint32_t processed_in_main = left_after_prologue & ~0x3;  // equivalent to (x / 4) * 4
     uint32_t left_after_main = left_after_prologue - processed_in_main;
 
     // epilogue
@@ -339,7 +339,7 @@ __device__ __forceinline__ void load_planar_tile(Tile tile[][kBlockSize / kStati
     tile[c][xy] = in.w;
   }
 
-  uint32_t processed_in_main = left_after_prologue & -4;  // equivalent to (x / 4) * 4
+  uint32_t processed_in_main = left_after_prologue & ~0x3;  // equivalent to (x / 4) * 4
   uint32_t left_after_main = left_after_prologue - processed_in_main;
 
   // epilogue
@@ -521,8 +521,6 @@ __device__ __forceinline__ void store_planar_hwc_pad(
   int64_t start_x = static_cast<int64_t>(blockIdx.x - sample.first_block) * kBlockSize;
   int64_t end_x = ::min(start_x + kBlockSize, sample.sample_size);
 
-  const auto *__restrict__ fill_values = static_cast<const float16 *>(sample.fill_values);
-
   // Preload the norm values so they are accessed via registers and not from gmem via pointer.
   Compute norm_mul[kOutChannels], norm_add[kOutChannels];
 
@@ -545,7 +543,6 @@ __device__ __forceinline__ void store_planar_hwc_pad(
 
   // TODO(klecki) in the version without mirror, we can keep one offset, as we can start the
   // output pointer at the output tile.
-  auto *out_aligned = sample.out;
   auto *out_h2 = reinterpret_cast<__half2 *>(sample.out);
   uint32_t to_write = end_x_padded - start_x_padded;
 
@@ -726,8 +723,6 @@ __global__ void Hwc2HwcNormalizePadFp16(const Hwc2HwcChwSampleDesc<Out, In> *sam
                                         uint32_t *first_blocks, uint32_t num_samples) {
   static_assert(std::is_same<In, uint8_t>::value, "Only uint8_t supported as input");
 
-  constexpr int kOutChannels = kStaticChannels + 1;
-
   int sample_idx = FindSampleIdx(first_blocks, num_samples);
   const auto sample = samples[sample_idx];
 

dali/operators/math/expressions/expression_impl_gpu_ternary.cuh

@@ -64,7 +64,6 @@ __global__ void ExecuteTiledTernaryOp1D(const SampleDescGPU<3> *samples, const T
   const auto &tile = tiles[blockIdx.y];
   const auto &sample = samples[tile.sample_idx];
   auto output = static_cast<Result *>(sample.output.data);
-  auto &out_strides = sample.output.strides;
   auto &arg0 = sample.args[0];
   auto &arg1 = sample.args[1];
   auto &arg2 = sample.args[2];