diff --git a/cmake/modules/CppLibrary.cmake b/cmake/modules/CppLibrary.cmake
index b0d04f9bc6..6d5b627a56 100644
--- a/cmake/modules/CppLibrary.cmake
+++ b/cmake/modules/CppLibrary.cmake
@@ -65,9 +65,6 @@ function(fbgemm_get_warning_flags)
     -Wmissing-braces
     -Wmismatched-tags
     -Waddress-of-packed-member
-    # The two flags below are on, but they do not stop the build. See the
-    # `-Wno-error=` lines in `_cc_suppressions_common`. Those lines give the
-    # condition to remove them.
     -Wshadow
     -Wzero-as-null-pointer-constant
     -Wunused-variable
@@ -178,16 +175,7 @@ function(fbgemm_get_warning_flags)
     -Wno-deprecated-enum-enum-conversion
     -Wno-strict-aliasing
     -Wno-sign-compare
-    -Wno-vla
-    -Wno-error=unused-parameter
-    -Wno-error=attributes
-    # Both compilers accept this line. Remove it when the warning count is
-    # zero.
-    -Wno-error=shadow
-    # Both compilers accept this line. The warning also occurs in
-    # third-party headers, and those headers need `-isystem`. Remove this line
-    # when the count is zero and the headers are system includes.
-    -Wno-error=zero-as-null-pointer-constant)
+    -Wno-vla)
 
   # Clang suppressions. The clang version controls which lines apply. The
   # CXX path uses the version of the host clang. The hipcc path uses all of
@@ -203,24 +191,16 @@ function(fbgemm_get_warning_flags)
     # Remove this line when the warning count is zero.
     -Wno-error=shorten-64-to-32)
 
-  set(_cc_suppressions_clang_gt13
-    -Wno-error=unused-but-set-parameter
-    -Wno-error=unused-but-set-variable)
-
   set(_cc_suppressions_clang_gt17
-    -Wno-vla-cxx-extension
-    -Wno-error=global-constructors)
+    -Wno-vla-cxx-extension)
 
   # Full clang-shaped suppression set, assembled unconditionally so it is
   # available even when the HOST compiler is GCC. Used only for the hipcc list.
   set(_cc_suppressions_clang
     ${_cc_suppressions_clang_base}
-    ${_cc_suppressions_clang_gt13}
     ${_cc_suppressions_clang_gt17})
 
   set(_cc_suppressions_gcc
-    -Wno-error=unused-but-set-parameter
-    -Wno-error=unused-but-set-variable
     -Wno-error=array-bounds
     -Wno-error=maybe-uninitialized)
 
@@ -231,10 +211,6 @@ function(fbgemm_get_warning_flags)
   if(CMAKE_CXX_COMPILER_ID MATCHES Clang)
     list(APPEND _cc_suppressions ${_cc_suppressions_clang_base})
 
-    if(CMAKE_CXX_COMPILER_VERSION VERSION_GREATER 13.0.0)
-      list(APPEND _cc_suppressions ${_cc_suppressions_clang_gt13})
-    endif()
-
     if(CMAKE_CXX_COMPILER_VERSION VERSION_GREATER 17.0.0)
       list(APPEND _cc_suppressions ${_cc_suppressions_clang_gt17})
     endif()
diff --git a/defs.bzl b/defs.bzl
index 83ca4ac357..b68ecf1ee3 100644
--- a/defs.bzl
+++ b/defs.bzl
@@ -29,12 +29,10 @@ def default_compiler_flags():
             # Buck exposes Folly as non-system headers, where its intentional
             # `__int128` use triggers the `-pedantic` diagnostic.
             "-Wno-error=pedantic",
-            # These four lines mirror the complete CMake GCC-specific list.
+            # These two lines mirror the complete CMake GCC-specific list.
             # Remove each line when its warning count is zero.
             "-Wno-error=array-bounds",
             "-Wno-error=maybe-uninitialized",
-            "-Wno-error=unused-but-set-parameter",
-            "-Wno-error=unused-but-set-variable",
             "-fno-trapping-math",
         ]
         + select({
diff --git a/fbgemm_gpu/codegen/training/optimizer/embedding_optimizer_split_device_kernel_template.cuh b/fbgemm_gpu/codegen/training/optimizer/embedding_optimizer_split_device_kernel_template.cuh
index 1d663529b4..3f7c777a7a 100644
--- a/fbgemm_gpu/codegen/training/optimizer/embedding_optimizer_split_device_kernel_template.cuh
+++ b/fbgemm_gpu/codegen/training/optimizer/embedding_optimizer_split_device_kernel_template.cuh
@@ -245,7 +245,7 @@ DEVICE_INLINE void {{ mdesc }}_{{ optimizer }}_table_update_kernel(
     const uint32_t run_id,
     const uint32_t cache_loc_run_id,
     const int32_t D,
-    const int32_t t,
+    [[maybe_unused]] const int32_t t,
     const int64_t idx,
     {%- if has_global_weight_decay_support %}
     const float global_weight_decay,
@@ -262,6 +262,10 @@ DEVICE_INLINE void {{ mdesc }}_{{ optimizer }}_table_update_kernel(
     {%- endfor %}
     {{ args.split_ref_kernel_args | replace_pta_namespace() | join(",\n    ") }}
 ) {
+    {%- for tensor in args.split_tensors %}
+    static_cast<void>({{ tensor }}_placements);
+    static_cast<void>({{ tensor }}_offsets);
+    {%- endfor %}
     constexpr auto kIsInt8 = std::is_same_v<emb_t, uint8_t>;
     // Copy value to max_vecs to make max_vecs_per_thread known at compile time
     // when kUseVecBlocking == false
diff --git a/fbgemm_gpu/experimental/gen_ai/src/attention/cuda/gqa_attn_splitk.cu b/fbgemm_gpu/experimental/gen_ai/src/attention/cuda/gqa_attn_splitk.cu
index 9a8238ece3..6b83831443 100644
--- a/fbgemm_gpu/experimental/gen_ai/src/attention/cuda/gqa_attn_splitk.cu
+++ b/fbgemm_gpu/experimental/gen_ai/src/attention/cuda/gqa_attn_splitk.cu
@@ -119,17 +119,20 @@ template <
     CacheLogicalDtype KVDataType>
 __global__ void __launch_bounds__(kThreadsPerWarp* kSplitKWarpsPerBlock, 1)
     gqa_attn_splitk_wmma_kernel(
-        const pta::
+        [[maybe_unused]] const pta::
             PackedTensorAccessor32<at::BFloat16, 4, at::RestrictPtrTraits> XQ,
-        const pta::PackedTensorAccessor64<kv_t, 4, at::RestrictPtrTraits>
-            cache_K,
-        const pta::PackedTensorAccessor64<kv_t, 4, at::RestrictPtrTraits>
-            cache_V,
-        pta::PackedTensorAccessor32<float, 4, at::RestrictPtrTraits> out_splitK,
-        const pta::PackedTensorAccessor32<int32_t, 1, at::RestrictPtrTraits>
-            seq_positions,
-        pta::PackedTensorAccessor32<float, 4, at::RestrictPtrTraits> metadata,
-        float qk_scale) {
+        [[maybe_unused]] const pta::
+            PackedTensorAccessor64<kv_t, 4, at::RestrictPtrTraits> cache_K,
+        [[maybe_unused]] const pta::
+            PackedTensorAccessor64<kv_t, 4, at::RestrictPtrTraits> cache_V,
+        [[maybe_unused]] pta::
+            PackedTensorAccessor32<float, 4, at::RestrictPtrTraits> out_splitK,
+        [[maybe_unused]] const pta::
+            PackedTensorAccessor32<int32_t, 1, at::RestrictPtrTraits>
+                seq_positions,
+        [[maybe_unused]] pta::
+            PackedTensorAccessor32<float, 4, at::RestrictPtrTraits> metadata,
+        [[maybe_unused]] float qk_scale) {
 #if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 800))
   // Need kWarpsPerBlock == blockDim.y;
   // Need D_H == 128
diff --git a/fbgemm_gpu/experimental/gen_ai/src/comm/car.cu b/fbgemm_gpu/experimental/gen_ai/src/comm/car.cu
index 9209138e85..82c9447572 100644
--- a/fbgemm_gpu/experimental/gen_ai/src/comm/car.cu
+++ b/fbgemm_gpu/experimental/gen_ai/src/comm/car.cu
@@ -96,7 +96,7 @@ __launch_bounds__(512)
 #endif
     __global__ void one_shot_all_reduce(
         int32_t rank,
-        int32_t world_size,
+        [[maybe_unused]] int32_t world_size,
         int32_t flag,
         std::array<int32_t*, 8> barriers,
         std::array<at::BFloat16*, 8> inputs,
@@ -212,7 +212,7 @@ __launch_bounds__(512)
 #endif
     __global__ void one_shot_all_reduce_pipelined(
         int32_t rank,
-        int32_t world_size,
+        [[maybe_unused]] int32_t world_size,
         int32_t flag,
         std::array<int32_t*, 8> barriers,
         std::array<at::BFloat16*, 8> inputs,
@@ -451,7 +451,7 @@ __launch_bounds__(512) __global__ void reduce_scatter(
 __launch_bounds__(1024) __global__ void reduce_scatter(
 #endif
     int32_t rank,
-    int32_t world_size,
+    [[maybe_unused]] int32_t world_size,
     int32_t flag,
     std::array<int32_t*, 8> barriers,
     std::array<at::BFloat16*, 8> inputs,
@@ -555,7 +555,7 @@ __launch_bounds__(512) __global__ void two_shot_all_reduce(
 __launch_bounds__(1024) __global__ void two_shot_all_reduce(
 #endif
     int32_t rank,
-    int32_t world_size,
+    [[maybe_unused]] int32_t world_size,
     int32_t flag,
     std::array<int32_t*, 8> barriers,
     std::array<at::BFloat16*, 8> inputs,
@@ -669,7 +669,7 @@ __launch_bounds__(512) __global__ void two_shot_all_reduce_pipelined(
 __launch_bounds__(1024) __global__ void two_shot_all_reduce_pipelined(
 #endif
     int32_t rank,
-    int32_t world_size,
+    [[maybe_unused]] int32_t world_size,
     int32_t flag,
     std::array<int32_t*, 8> barriers,
     std::array<at::BFloat16*, 8> inputs,
@@ -785,7 +785,7 @@ void one_shot_car_allreduce(
     at::Tensor y_allreduce,
     at::Tensor y,
     std::optional<at::Tensor> z,
-    int64_t comm_idx,
+    [[maybe_unused]] int64_t comm_idx,
     bool enable_pipelining) { // match the API with nccl_allreduce in
                               // https://fburl.com/code/v538vig9
   c10::cuda::CUDAGuard gg(y_allreduce.device());
@@ -897,7 +897,7 @@ void two_shot_car_allreduce(
     at::Tensor y_allreduce,
     at::Tensor y,
     std::optional<at::Tensor> z,
-    int64_t comm_idx,
+    [[maybe_unused]] int64_t comm_idx,
     bool enable_pipelining) { // match the API with nccl_allreduce in
                               // https://fburl.com/code/v538vig9
   c10::cuda::CUDAGuard gg(y_allreduce.device());
@@ -1000,8 +1000,8 @@ void car_reducescatter(
     at::Tensor dst,
     at::Tensor src,
     bool split_last_dim,
-    int64_t comm_idx) { // match the API with nccl_allreduce in
-                        // https://fburl.com/code/v538vig9
+    [[maybe_unused]] int64_t comm_idx) { // match the API with nccl_allreduce in
+  // https://fburl.com/code/v538vig9
   auto state = get_car_state();
   c10::cuda::CUDAGuard gg(dst.device());
   TORCH_CHECK(dst.is_contiguous());
diff --git a/fbgemm_gpu/experimental/gen_ai/src/kv_cache/kv_cache.cpp b/fbgemm_gpu/experimental/gen_ai/src/kv_cache/kv_cache.cpp
index c4d644a98f..4aae575282 100644
--- a/fbgemm_gpu/experimental/gen_ai/src/kv_cache/kv_cache.cpp
+++ b/fbgemm_gpu/experimental/gen_ai/src/kv_cache/kv_cache.cpp
@@ -253,7 +253,7 @@ at::Tensor quantize_qkv_per_head_meta(
     at::Tensor /* varseq_seqpos */,
     std::optional<at::Tensor> /* varseq_batch */,
     std::optional<at::Tensor> /* is_precalculated_qparam */,
-    at::Tensor cache_K /* cache_K */,
+    at::Tensor /* cache_K */,
     at::Tensor /* cache_V */,
     at::Tensor /* XQ_O */,
     int64_t /* B */,
@@ -263,10 +263,10 @@ at::Tensor quantize_qkv_per_head_meta(
 }
 
 void convert_e4m3fn_kv_cache_to_e4m3fnuz_inplace_meta(
-    at::Tensor cache_K,
-    at::Tensor cache_V,
-    at::Tensor qparam_K,
-    at::Tensor qparam_v) {};
+    at::Tensor /* cache_K */,
+    at::Tensor /* cache_V */,
+    at::Tensor /* qparam_K */,
+    at::Tensor /* qparam_v */){};
 
 TORCH_LIBRARY_IMPL(fbgemm, Meta, m) {
   m.impl("rope_qkv_varseq_prefill", rope_qkv_varseq_prefill_meta);
diff --git a/fbgemm_gpu/experimental/gen_ai/src/kv_cache/kv_cache_dequantize.cu b/fbgemm_gpu/experimental/gen_ai/src/kv_cache/kv_cache_dequantize.cu
index 281c9509b5..223ac471e8 100644
--- a/fbgemm_gpu/experimental/gen_ai/src/kv_cache/kv_cache_dequantize.cu
+++ b/fbgemm_gpu/experimental/gen_ai/src/kv_cache/kv_cache_dequantize.cu
@@ -296,20 +296,30 @@ __global__ void dequantize_fp8_cache_kernel(
 
 __global__ void dequantize_fp8_cache_kernel_paged(
     // This code currently represents FP8 version not int4
-    pta::PackedTensorAccessor64<uint8_t, 4, at::RestrictPtrTraits>
-        cache_K, // [1][MAX_PAGE * PAGE_SIZE][N_KVH][D_H]
-    pta::PackedTensorAccessor64<uint8_t, 4, at::RestrictPtrTraits>
+    [[maybe_unused]] pta::PackedTensorAccessor64<
+        uint8_t,
+        4,
+        at::RestrictPtrTraits> cache_K, // [1][MAX_PAGE * PAGE_SIZE][N_KVH][D_H]
+    [[maybe_unused]] pta::PackedTensorAccessor64<
+        uint8_t,
+        4,
+        at::RestrictPtrTraits>
         cache_V, // [1][MAX_PAGE * PAGE_SIZE][N_KVH][D_H // G]
-    pta::PackedTensorAccessor32<int32_t, 1, at::RestrictPtrTraits> kv_seqlen,
-    pta::PackedTensorAccessor64<at::BFloat16, 4, at::RestrictPtrTraits>
-        cache_K_dq, // [1][MAX_T][N_KVH][D_H]
-    pta::PackedTensorAccessor64<at::BFloat16, 4, at::RestrictPtrTraits>
-        cache_V_dq, // [1][MAX_T][N_KVH][D_H]
-    int32_t* qparam_k_ptr,
-    int32_t* qparam_v_ptr,
-    int32_t* block_tables,
-    int32_t block_tables_b_stride,
-    int32_t page_size) {
+    [[maybe_unused]] pta::
+        PackedTensorAccessor32<int32_t, 1, at::RestrictPtrTraits> kv_seqlen,
+    [[maybe_unused]] pta::PackedTensorAccessor64<
+        at::BFloat16,
+        4,
+        at::RestrictPtrTraits> cache_K_dq, // [1][MAX_T][N_KVH][D_H]
+    [[maybe_unused]] pta::PackedTensorAccessor64<
+        at::BFloat16,
+        4,
+        at::RestrictPtrTraits> cache_V_dq, // [1][MAX_T][N_KVH][D_H]
+    [[maybe_unused]] int32_t* qparam_k_ptr,
+    [[maybe_unused]] int32_t* qparam_v_ptr,
+    [[maybe_unused]] int32_t* block_tables,
+    [[maybe_unused]] int32_t block_tables_b_stride,
+    [[maybe_unused]] int32_t page_size) {
   CUDA_KERNEL_ASSERT(0 && "unimplemented");
 }
 
diff --git a/fbgemm_gpu/experimental/gen_ai/src/moe/index_shuffling.cpp b/fbgemm_gpu/experimental/gen_ai/src/moe/index_shuffling.cpp
index 537ef56894..b7e0feac27 100644
--- a/fbgemm_gpu/experimental/gen_ai/src/moe/index_shuffling.cpp
+++ b/fbgemm_gpu/experimental/gen_ai/src/moe/index_shuffling.cpp
@@ -46,9 +46,9 @@ std::tuple<at::Tensor, at::Tensor, at::Tensor> index_shuffling_torch(
 
 std::tuple<at::Tensor, at::Tensor, at::Tensor> index_shuffling_torch_meta(
     const at::Tensor& routing_scores,
-    const std::optional<int64_t>& expert_index_start,
-    const std::optional<int64_t>& expert_index_end,
-    const std::optional<at::Tensor>& valid_token_count,
+    const std::optional<int64_t>& /* expert_index_start */,
+    const std::optional<int64_t>& /* expert_index_end */,
+    const std::optional<at::Tensor>& /* valid_token_count */,
     const int64_t top_k = 1) {
   auto T = routing_scores.sym_size(0);
   auto E = routing_scores.sym_size(1);
diff --git a/fbgemm_gpu/experimental/gen_ai/src/quantize/ck_extensions/bf16_grouped/bf16_grouped_gemm.hip b/fbgemm_gpu/experimental/gen_ai/src/quantize/ck_extensions/bf16_grouped/bf16_grouped_gemm.hip
index d08757cb04..3df3d08f7a 100644
--- a/fbgemm_gpu/experimental/gen_ai/src/quantize/ck_extensions/bf16_grouped/bf16_grouped_gemm.hip
+++ b/fbgemm_gpu/experimental/gen_ai/src/quantize/ck_extensions/bf16_grouped/bf16_grouped_gemm.hip
@@ -285,7 +285,7 @@ __global__ void set_kernel_args_m_sizes_kernel(
     BDataType* B,
     CDataType* output,
     int64_t* M_sizes,
-    int64_t M,
+    [[maybe_unused]] int64_t M,
     int64_t N,
     int64_t K,
     int64_t group_count) {
@@ -520,8 +520,8 @@ at::Tensor bf16bf16bf16_grouped_stacked(
     at::Tensor X,
     at::Tensor W,
     at::Tensor M_sizes,
-    std::optional<at::Tensor> out,
-    std::optional<int64_t> num_sms) {
+    [[maybe_unused]] std::optional<at::Tensor> out,
+    [[maybe_unused]] std::optional<int64_t> num_sms) {
   // Check that input datatypes are valid.
   // First confirm that there are the same number of groups in all inputs.
   int64_t group_count = M_sizes.size(0);
diff --git a/fbgemm_gpu/experimental/gen_ai/src/quantize/ck_extensions/fp8_rowwise_batched/fp8_rowwise_batched_gemm.hip b/fbgemm_gpu/experimental/gen_ai/src/quantize/ck_extensions/fp8_rowwise_batched/fp8_rowwise_batched_gemm.hip
index 299d667ced..9388d16fe9 100644
--- a/fbgemm_gpu/experimental/gen_ai/src/quantize/ck_extensions/fp8_rowwise_batched/fp8_rowwise_batched_gemm.hip
+++ b/fbgemm_gpu/experimental/gen_ai/src/quantize/ck_extensions/fp8_rowwise_batched/fp8_rowwise_batched_gemm.hip
@@ -202,7 +202,7 @@ at::Tensor f8f8bf16_rowwise_batched(
     at::Tensor WQ,
     at::Tensor x_scale,
     at::Tensor w_scale,
-    std::optional<at::Tensor> bias,
+    [[maybe_unused]] std::optional<at::Tensor> bias,
     bool use_fast_accum,
     std::optional<at::Tensor> output = std::nullopt) {
   // Check that input datatypes are valid.
diff --git a/fbgemm_gpu/experimental/gen_ai/src/quantize/quantize.cpp b/fbgemm_gpu/experimental/gen_ai/src/quantize/quantize.cpp
index a8d78b859f..ec91bce728 100644
--- a/fbgemm_gpu/experimental/gen_ai/src/quantize/quantize.cpp
+++ b/fbgemm_gpu/experimental/gen_ai/src/quantize/quantize.cpp
@@ -392,8 +392,8 @@ TORCH_LIBRARY_IMPL(fbgemm, CPU, m) {
 at::Tensor i8i8bf16_meta(
     at::Tensor XQ, // INT8
     at::Tensor WQ, // INT8
-    double scale,
-    int64_t split_k) {
+    [[maybe_unused]] double scale,
+    [[maybe_unused]] int64_t split_k) {
   const at::SymInt M = XQ.sym_size(0);
   const at::SymInt N = WQ.sym_size(0);
   auto Y = at::empty_symint({M, N}, XQ.options().dtype(at::kBFloat16));
@@ -539,8 +539,8 @@ at::Tensor f8f8bf16_cublas_meta(
 at::Tensor f8f8bf16_meta(
     at::Tensor X,
     at::Tensor W,
-    at::Tensor scale,
-    bool use_fast_accum = true) {
+    [[maybe_unused]] at::Tensor scale,
+    [[maybe_unused]] bool use_fast_accum = true) {
   const at::SymInt M = X.sym_size(0);
   const at::SymInt N = W.sym_size(0);
   auto Y = at::empty_symint({M, N}, X.options().dtype(at::kBFloat16));
@@ -608,7 +608,10 @@ at::Tensor f8f8bf16_tensorwise_meta(
   return Y;
 }
 
-at::Tensor f8f8bf16_lite_meta(at::Tensor X, at::Tensor W, at::Tensor scale) {
+at::Tensor f8f8bf16_lite_meta(
+    at::Tensor X,
+    at::Tensor W,
+    [[maybe_unused]] at::Tensor scale) {
   const at::SymInt M = X.sym_size(0);
   const at::SymInt N = W.sym_size(0);
   auto Y = at::empty_symint({M, N}, X.options().dtype(at::kBFloat16));
@@ -773,10 +776,10 @@ std::vector<at::Tensor> quantize_fp8_per_row_meta(
 }
 
 void scaled_fp4_quant_meta(
-    at::Tensor const& output,
-    at::Tensor const& input,
-    at::Tensor const& output_sf,
-    at::Tensor const& input_sf) {
+    [[maybe_unused]] at::Tensor const& output,
+    [[maybe_unused]] at::Tensor const& input,
+    [[maybe_unused]] at::Tensor const& output_sf,
+    [[maybe_unused]] at::Tensor const& input_sf) {
   return;
 }
 
diff --git a/fbgemm_gpu/experimental/gen_ai/src/quantize/quantize.cu b/fbgemm_gpu/experimental/gen_ai/src/quantize/quantize.cu
index 7ff5fd0d65..7c195a975f 100644
--- a/fbgemm_gpu/experimental/gen_ai/src/quantize/quantize.cu
+++ b/fbgemm_gpu/experimental/gen_ai/src/quantize/quantize.cu
@@ -439,7 +439,7 @@ __global__ void scaleMatrix1(
     T_S const* const input_scale,
     T_IN const* const input,
     const int64_t numel,
-    const int64_t lda) {
+    [[maybe_unused]] const int64_t lda) {
   for (int64_t i = threadIdx.x + blockIdx.x * blockDim.x; i < numel;
        i += (size_t)blockDim.x * gridDim.x) {
     output[i] = T_OUT(
@@ -454,7 +454,7 @@ __global__ void scaleMatrix2(
     T_S const* const input_scale,
     T_IN const* const input,
     const int64_t numel,
-    const int64_t lda,
+    [[maybe_unused]] const int64_t lda,
     at::PhiloxCudaState stochastic_rounding_philox_args) {
   auto stoc_rounding_state = StochasticRoundingRNGState(
       stochastic_rounding_philox_args, threadIdx.x + blockIdx.x * blockDim.x);
@@ -703,7 +703,7 @@ __global__ void computeFP8QuantizeScale(
     T_S* const quant_ptr,
     const T_W* const weights,
     const int64_t size,
-    const int64_t n,
+    [[maybe_unused]] const int64_t n,
     const int64_t total_elements_per_slice,
     const int64_t* bs,
     const float* scale_ub) {
@@ -837,7 +837,7 @@ at::Tensor get_fp8_per_tensor_scale(
 at::Tensor quantize_fp8_per_tensor_fixed_scale(
     at::Tensor input,
     at::Tensor scale,
-    std::optional<at::Tensor> bs, // batch size
+    [[maybe_unused]] std::optional<at::Tensor> bs, // batch size
     bool stochastic_rounding) {
   CUDA_DEVICE_GUARD(input);
   TORCH_CHECK(
@@ -1308,10 +1308,10 @@ __global__ void fused_quantize_rowwise(
 
 std::vector<at::Tensor> quantize_fp8_per_row(
     at::Tensor input,
-    std::optional<at::Tensor> bs, // batch size
+    [[maybe_unused]] std::optional<at::Tensor> bs, // batch size
     std::optional<at::Tensor> scale_ub, // scale upperbound
     std::optional<c10::ScalarType> output_dtype, // Quantization type
-    bool stochastic_rounding) {
+    [[maybe_unused]] bool stochastic_rounding) {
   TORCH_CHECK(input.dim() >= 2, "Invalid dim. The dim of input should be >= 2");
   TORCH_CHECK(
       input.scalar_type() == torch::kBFloat16 ||
@@ -1381,8 +1381,8 @@ std::vector<at::Tensor> quantize_fp8_per_row(
 
 std::vector<at::Tensor> quantize_fp8_per_col(
     at::Tensor input,
-    std::optional<at::Tensor> bs, // batch size
-    std::optional<at::Tensor> scale_ub) // scale upperbound)
+    [[maybe_unused]] std::optional<at::Tensor> bs, // batch size
+    [[maybe_unused]] std::optional<at::Tensor> scale_ub) // scale upperbound)
 {
   CUDA_DEVICE_GUARD(input);
   TORCH_CHECK(
@@ -2241,18 +2241,19 @@ void scaled_fp4_quant(
 }
 #else
 void scaled_fp4_quant(
-    at::Tensor const& output,
-    at::Tensor const& input,
-    at::Tensor const& output_sf,
-    at::Tensor const& input_sf) {
+    [[maybe_unused]] at::Tensor const& output,
+    [[maybe_unused]] at::Tensor const& input,
+    [[maybe_unused]] at::Tensor const& output_sf,
+    [[maybe_unused]] at::Tensor const& input_sf) {
   throw std::runtime_error(
       "CUDA version is older than 12.8"); // requires CUDA>=12.8
 }
 std::vector<at::Tensor> fake_quantize_nvfp4_per_tensor(
-    at::Tensor input,
-    std::optional<at::Tensor> static_scales, // static scale is optional
-    std::optional<at::Tensor> bs, // batch size
-    std::optional<at::Tensor> scale_ub) // scale upperbound)
+    [[maybe_unused]] at::Tensor input,
+    [[maybe_unused]] std::optional<at::Tensor>
+        static_scales, // static scale is optional
+    [[maybe_unused]] std::optional<at::Tensor> bs, // batch size
+    [[maybe_unused]] std::optional<at::Tensor> scale_ub) // scale upperbound)
 {
   throw std::runtime_error(
       "CUDA version is older than 12.8"); // requires CUDA>=12.8
@@ -2261,62 +2262,63 @@ std::vector<at::Tensor> fake_quantize_nvfp4_per_tensor(
 
 #else
 std::vector<at::Tensor> quantize_fp8_per_tensor(
-    at::Tensor input,
-    std::optional<at::Tensor> bs, // batch size
-    std::optional<at::Tensor> scale_ub,
-    bool stochastic_rounding) { // scale upperbound
+    [[maybe_unused]] at::Tensor input,
+    [[maybe_unused]] std::optional<at::Tensor> bs, // batch size
+    [[maybe_unused]] std::optional<at::Tensor> scale_ub,
+    [[maybe_unused]] bool stochastic_rounding) { // scale upperbound
   throw std::runtime_error(
       "CUDA version is older than 12.0"); // requires CUDA>=12
 }
 
 std::vector<at::Tensor> quantize_fp8_per_row(
-    at::Tensor input,
-    std::optional<at::Tensor> bs, // batch size
-    std::optional<at::Tensor> scale_ub, // scale upperbound
-    std::optional<c10::ScalarType> output_dtype,
-    bool stochastic_rounding) { // quantization type
+    [[maybe_unused]] at::Tensor input,
+    [[maybe_unused]] std::optional<at::Tensor> bs, // batch size
+    [[maybe_unused]] std::optional<at::Tensor> scale_ub, // scale upperbound
+    [[maybe_unused]] std::optional<c10::ScalarType> output_dtype,
+    [[maybe_unused]] bool stochastic_rounding) { // quantization type
   throw std::runtime_error(
       "CUDA version is older than 12.0"); // requires CUDA>=12
 }
 
 at::Tensor quantize_fp8_per_tensor_fixed_scale(
-    at::Tensor input,
-    at::Tensor scale,
-    std::optional<at::Tensor> bs,
-    bool stochastic_rounding) { // batch size
+    [[maybe_unused]] at::Tensor input,
+    [[maybe_unused]] at::Tensor scale,
+    [[maybe_unused]] std::optional<at::Tensor> bs,
+    [[maybe_unused]] bool stochastic_rounding) { // batch size
   throw std::runtime_error(
       "CUDA version is older than 12.0"); // requires CUDA>=12
 }
 
 at::Tensor get_fp8_per_tensor_scale(
-    at::Tensor input,
-    std::optional<at::Tensor> bs, // batch size
-    std::optional<at::Tensor> scale_ub) { // scale upperbound
+    [[maybe_unused]] at::Tensor input,
+    [[maybe_unused]] std::optional<at::Tensor> bs, // batch size
+    [[maybe_unused]] std::optional<at::Tensor> scale_ub) { // scale upperbound
   throw std::runtime_error(
       "CUDA version is older than 12.0"); // requires CUDA>=12
 }
 
 std::vector<at::Tensor> quantize_fp8_per_col(
-    at::Tensor input,
-    std::optional<at::Tensor> bs, // batch size
-    std::optional<at::Tensor> scale_ub) { // scale upperbound
+    [[maybe_unused]] at::Tensor input,
+    [[maybe_unused]] std::optional<at::Tensor> bs, // batch size
+    [[maybe_unused]] std::optional<at::Tensor> scale_ub) { // scale upperbound
   throw std::runtime_error(
       "CUDA version is older than 12.0"); // requires CUDA>=12
 }
 
 void scaled_fp4_quant(
-    at::Tensor const& output,
-    at::Tensor const& input,
-    at::Tensor const& output_sf,
-    at::Tensor const& input_sf) {
+    [[maybe_unused]] at::Tensor const& output,
+    [[maybe_unused]] at::Tensor const& input,
+    [[maybe_unused]] at::Tensor const& output_sf,
+    [[maybe_unused]] at::Tensor const& input_sf) {
   throw std::runtime_error(
       "CUDA version is older than 12.8"); // requires CUDA>=12.8
 }
 std::vector<at::Tensor> fake_quantize_nvfp4_per_tensor(
-    at::Tensor input,
-    std::optional<at::Tensor> static_scales, // static scale is optional
-    std::optional<at::Tensor> bs, // batch size
-    std::optional<at::Tensor> scale_ub) // scale upperbound)
+    [[maybe_unused]] at::Tensor input,
+    [[maybe_unused]] std::optional<at::Tensor>
+        static_scales, // static scale is optional
+    [[maybe_unused]] std::optional<at::Tensor> bs, // batch size
+    [[maybe_unused]] std::optional<at::Tensor> scale_ub) // scale upperbound)
 {
   throw std::runtime_error(
       "CUDA version is older than 12.8"); // requires CUDA>=12.8
diff --git a/fbgemm_gpu/include/fbgemm_gpu/embedding_common.h b/fbgemm_gpu/include/fbgemm_gpu/embedding_common.h
index 860aa18f43..ccd001bc1e 100644
--- a/fbgemm_gpu/include/fbgemm_gpu/embedding_common.h
+++ b/fbgemm_gpu/include/fbgemm_gpu/embedding_common.h
@@ -84,7 +84,7 @@ enum class BoundsCheckMode : uint8_t {
 // what device kernels (whose format is selected per-arch at device-compile
 // time) read and write.
 inline at::ScalarType getNFP8ScalarType(
-    const c10::DeviceIndex device_index = -1) {
+    [[maybe_unused]] const c10::DeviceIndex device_index = -1) {
 #ifdef USE_ROCM
   // fnuz archs: the gfx94x family (gfx940/941/942, MI300) and gfx90a. The
   // substring match mirrors split_embedding_configs.py:nfp8_dtype; keep the
diff --git a/fbgemm_gpu/include/fbgemm_gpu/rocm/split_embeddings_common.h b/fbgemm_gpu/include/fbgemm_gpu/rocm/split_embeddings_common.h
index 2065e590e7..584750c621 100644
--- a/fbgemm_gpu/include/fbgemm_gpu/rocm/split_embeddings_common.h
+++ b/fbgemm_gpu/include/fbgemm_gpu/rocm/split_embeddings_common.h
@@ -148,10 +148,10 @@ __device__ void llvm_amdgcn_raw_buffer_store_fp32x2(
 template <typename emb_t, int32_t embedding_dim, typename index_t>
 struct load_row_per_warp {
   static __device__ void run(
-      emb_t* emb_data,
-      index_t row_index,
-      const emb_t* p_emb_table,
-      int lane_id) {
+      emb_t* /* emb_data */,
+      index_t /* row_index */,
+      const emb_t* /* p_emb_table */,
+      int /* lane_id */) {
     // Types are not supported, but we need an instance of run method to avoid
     // run-time .so symbol failure. Currently, the kernel dispatch for
     // unsupported type is guarded on host side
@@ -255,8 +255,11 @@ template <
 struct accumulate_row_per_warp {
   static constexpr int dword_per_row =
       (embedding_dim + THREADS_PER_ROW - 1) / THREADS_PER_ROW;
-  static __device__ void
-  run(output_t* acc, emb_t* emb_data, int lane_id, float row_weight = 1.0) {
+  static __device__ void run(
+      output_t* acc,
+      emb_t* emb_data,
+      int /* lane_id */,
+      float row_weight = 1.0) {
     if constexpr (!weighted) {
 #pragma unroll
       for (int i = 0; i < dword_per_row; i++) {
@@ -279,7 +282,8 @@ struct accumulate_row_per_warp {
 
 template <typename emb_t, int32_t embedding_dim>
 struct store_row_per_warp {
-  static __device__ void run(const emb_t* acc, emb_t* p_output, int lane_id) {
+  static __device__ void
+  run(const emb_t* /* acc */, emb_t* /* p_output */, int /* lane_id */) {
     // Types are not supported, but we need an instance of run method to avoid
     // run-time .so symbol failure. Currently, the kernel dispatch for
     // unsupported type is guarded on host function
@@ -500,7 +504,7 @@ __device__ __forceinline__ void generic_dpp_reduction(data_t& result) {
 // Use corresponding assebly instruction for dpp reduction in case
 // of trivial operation with an option to use custom operation
 template <typename data_t, typename reduce_op_t, int wave_size = 64>
-__device__ __forceinline__ void dpp_reduction(data_t& result) {
+__device__ __forceinline__ void dpp_reduction([[maybe_unused]] data_t& result) {
 #if defined(__gfx942__) || defined(__gfx90a__) || defined(__gfx950__)
   if constexpr (std::is_same_v<reduce_op_t, reduce_op::sum>) {
     DPP_REDUCE_F16_F32(add);
diff --git a/fbgemm_gpu/include/fbgemm_gpu/utils/cuda_prelude.cuh b/fbgemm_gpu/include/fbgemm_gpu/utils/cuda_prelude.cuh
index 17f5083539..c8647bedd4 100755
--- a/fbgemm_gpu/include/fbgemm_gpu/utils/cuda_prelude.cuh
+++ b/fbgemm_gpu/include/fbgemm_gpu/utils/cuda_prelude.cuh
@@ -94,7 +94,8 @@ DEVICE_INLINE T shfl_xor(
     const T val,
     int laneMask,
     int width = kWarpSize,
-    unsigned shfl_sync_mask = static_cast<unsigned>(kFullWarpMask)) {
+    [[maybe_unused]] unsigned shfl_sync_mask =
+        static_cast<unsigned>(kFullWarpMask)) {
 #if defined(USE_ROCM)
   return __shfl_xor(val, laneMask, width);
 #else
@@ -107,7 +108,8 @@ DEVICE_INLINE T shfl_sync(
     const T val,
     int srcLane = 0,
     int width = kWarpSize,
-    unsigned shfl_sync_mask = static_cast<unsigned>(kFullWarpMask)) {
+    [[maybe_unused]] unsigned shfl_sync_mask =
+        static_cast<unsigned>(kFullWarpMask)) {
 #if defined(USE_ROCM)
   return __shfl(val, srcLane, width);
 #else
@@ -120,7 +122,8 @@ DEVICE_INLINE T shfl_down_sync(
     const T val,
     unsigned delta,
     int width = kWarpSize,
-    unsigned shfl_sync_mask = static_cast<unsigned>(kFullWarpMask)) {
+    [[maybe_unused]] unsigned shfl_sync_mask =
+        static_cast<unsigned>(kFullWarpMask)) {
 #if defined(USE_ROCM)
   return __shfl_down(val, delta, width);
 #else
@@ -134,7 +137,8 @@ DEVICE_INLINE uint64_t ballot_sync(
 DEVICE_INLINE uint32_t ballot_sync(
 #endif
     int predicate,
-    unsigned shfl_sync_mask = static_cast<unsigned>(kFullWarpMask)) {
+    [[maybe_unused]] unsigned shfl_sync_mask =
+        static_cast<unsigned>(kFullWarpMask)) {
 #if defined(USE_ROCM)
   return __ballot(predicate);
 #else
@@ -156,7 +160,8 @@ DEVICE_INLINE uint32_t ballot_sync(
 template <typename T, int ReduceWidth = kWarpSize>
 DEVICE_INLINE T warpReduceAllSum(
     T val,
-    unsigned shfl_sync_mask = static_cast<unsigned>(kFullWarpMask)) {
+    [[maybe_unused]] unsigned shfl_sync_mask =
+        static_cast<unsigned>(kFullWarpMask)) {
 #if defined(USE_ROCM) && \
     (defined(__gfx942__) || defined(__gfx90a__) || defined(__gfx950__))
   return rocm::wave_reduce<
diff --git a/fbgemm_gpu/include/fbgemm_gpu/utils/stochastic_rounding.cuh b/fbgemm_gpu/include/fbgemm_gpu/utils/stochastic_rounding.cuh
index 5fff80de0d..a5211268b9 100644
--- a/fbgemm_gpu/include/fbgemm_gpu/utils/stochastic_rounding.cuh
+++ b/fbgemm_gpu/include/fbgemm_gpu/utils/stochastic_rounding.cuh
@@ -122,7 +122,7 @@ template <typename dst_t, typename src_t>
 DEVICE_INLINE void stochastic_rounding_vector(
     dst_t* output,
     const Vec4T<src_t>& value,
-    StochasticRoundingRNGState& state,
+    StochasticRoundingRNGState& /* state */,
     const float2 /* not used */) {
   value.store(output);
 }
diff --git a/fbgemm_gpu/include/fbgemm_gpu/utils/vec4.cuh b/fbgemm_gpu/include/fbgemm_gpu/utils/vec4.cuh
index 37887d12f1..4a44b89061 100644
--- a/fbgemm_gpu/include/fbgemm_gpu/utils/vec4.cuh
+++ b/fbgemm_gpu/include/fbgemm_gpu/utils/vec4.cuh
@@ -142,8 +142,8 @@ struct Vec4T<float> : public Vec4BaseT<float> {
     acc = static_cast<float4>(fp8_ptr[0]);
   }
 
-  DEVICE_INLINE void load(const uint8_t* p) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void load(const uint8_t* /* p */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
   DEVICE_INLINE void store(float* p) const {
@@ -198,8 +198,8 @@ struct Vec4T<float> : public Vec4BaseT<float> {
     fp8_ptr[0] = static_cast<__nv_fp8x4_e4m3>(acc);
   }
 
-  DEVICE_INLINE void store(uint8_t* p) const {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void store(uint8_t* /* p */) const {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
   DEVICE_INLINE static void copy(const float* src, float* dst) {
@@ -322,8 +322,8 @@ struct Vec4T<at::Half> : public Vec4BaseT<at::Half> {
     acc = *((const float4*)p);
   }
 
-  DEVICE_INLINE void load(const uint8_t* p) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void load(const uint8_t* /* p */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
   DEVICE_INLINE void store(at::Half* p) const {
@@ -364,8 +364,8 @@ struct Vec4T<at::Half> : public Vec4BaseT<at::Half> {
     *((float4*)p) = acc;
   }
 
-  DEVICE_INLINE void store(uint8_t* p) const {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void store(uint8_t* /* p */) const {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
   DEVICE_INLINE static void copy(const at::Half* src, at::Half* dst) {
@@ -521,8 +521,8 @@ struct Vec4T<at::BFloat16> : public Vec4BaseT<at::BFloat16> {
     acc = *((const float4*)p);
   }
 
-  DEVICE_INLINE void load(const uint8_t* p) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void load(const uint8_t* /* p */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
   DEVICE_INLINE void store(at::Half* p) const {
@@ -563,8 +563,8 @@ struct Vec4T<at::BFloat16> : public Vec4BaseT<at::BFloat16> {
     *((float4*)p) = acc;
   }
 
-  DEVICE_INLINE void store(uint8_t* p) const {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void store(uint8_t* /* p */) const {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
   DEVICE_INLINE static void copy(const at::BFloat16* src, at::BFloat16* dst) {
diff --git a/fbgemm_gpu/include/fbgemm_gpu/utils/vec4acc.cuh b/fbgemm_gpu/include/fbgemm_gpu/utils/vec4acc.cuh
index a9e96cadc7..7e1472ed17 100644
--- a/fbgemm_gpu/include/fbgemm_gpu/utils/vec4acc.cuh
+++ b/fbgemm_gpu/include/fbgemm_gpu/utils/vec4acc.cuh
@@ -92,8 +92,8 @@ struct Vec4AccT {
   }
 #endif
 
-  DEVICE_INLINE void store(uint8_t* ptr) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void store(uint8_t* /* ptr */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
   DEVICE_INLINE void add(const float4* ptr) {
@@ -120,28 +120,32 @@ struct Vec4AccT {
     this->fma_(vals, weight);
   }
 
-  DEVICE_INLINE void add(const uint8_t* ptr) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void add(const uint8_t* /* ptr */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void fma(const uint8_t* ptr, const float weight) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void fma(const uint8_t* /* ptr */, const float /* weight */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void add(const c10::Float8_e4m3fn* ptr) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void add(const c10::Float8_e4m3fn* /* ptr */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void fma(const c10::Float8_e4m3fn* ptr, const float weight) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void fma(
+      const c10::Float8_e4m3fn* /* ptr */,
+      const float /* weight */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void add(const c10::Float8_e4m3fnuz* ptr) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void add(const c10::Float8_e4m3fnuz* /* ptr */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void fma(const c10::Float8_e4m3fnuz* ptr, const float weight) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void fma(
+      const c10::Float8_e4m3fnuz* /* ptr */,
+      const float /* weight */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
   DEVICE_INLINE void div(uint32_t denom) {
@@ -208,8 +212,8 @@ struct Vec4StepT<STEP, float> : Vec4AccT {
   }
 #endif
 
-  DEVICE_INLINE void index_store(uint32_t idx, uint8_t* ptr) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_store(uint32_t /* idx */, uint8_t* /* ptr */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
   DEVICE_INLINE void
@@ -246,9 +250,11 @@ struct Vec4StepT<STEP, float> : Vec4AccT {
   }
 #endif
 
-  DEVICE_INLINE void
-  index_weighted_store(uint32_t idx, uint8_t* ptr, const float weight) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_weighted_store(
+      uint32_t /* idx */,
+      uint8_t* /* ptr */,
+      const float /* weight */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 };
 
@@ -340,8 +346,8 @@ struct Vec4StepT<STEP, at::Half> : Vec4AccT {
   }
 #endif
 
-  DEVICE_INLINE void index_store(uint32_t idx, uint8_t* ptr) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_store(uint32_t /* idx */, uint8_t* /* ptr */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
   DEVICE_INLINE void
@@ -378,76 +384,86 @@ struct Vec4StepT<STEP, at::Half> : Vec4AccT {
   }
 #endif
 
-  DEVICE_INLINE void
-  index_weighted_store(uint32_t idx, uint8_t* ptr, const float weight) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_weighted_store(
+      uint32_t /* idx */,
+      uint8_t* /* ptr */,
+      const float /* weight */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 };
 
 template <uint32_t STEP>
 struct Vec4StepT<STEP, uint8_t> : Vec4AccT {
   DEVICE_INLINE Vec4StepT() {
-    CUDA_KERNEL_ASSERT(false);
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void load(const uint8_t* ptr, const uint32_t idx) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void load(const uint8_t* /* ptr */, const uint32_t /* idx */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
   DEVICE_INLINE void sum() {
-    CUDA_KERNEL_ASSERT(false);
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
   DEVICE_INLINE void weighted_sum(
-      const float* const weights,
-      const uint32_t idx_shift,
-      const uint32_t idx_scale) {
-    CUDA_KERNEL_ASSERT(false);
+      const float* const /* weights */,
+      const uint32_t /* idx_shift */,
+      const uint32_t /* idx_scale */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void index_add(uint32_t idx) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_add(uint32_t /* idx */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void index_fma(uint32_t idx, const float weight) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_fma(uint32_t /* idx */, const float /* weight */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void index_store(uint32_t idx, float4* ptr) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_store(uint32_t /* idx */, float4* /* ptr */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void index_store(uint32_t idx, float2* ptr) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_store(uint32_t /* idx */, float2* /* ptr */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void index_store(uint32_t idx, uint8_t* ptr) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_store(uint32_t /* idx */, uint8_t* /* ptr */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void
-  index_weighted_store(uint32_t idx, float4* ptr, const float weight) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_weighted_store(
+      uint32_t /* idx */,
+      float4* /* ptr */,
+      const float /* weight */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void
-  index_weighted_store(uint32_t idx, float2* ptr, const float weight) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_weighted_store(
+      uint32_t /* idx */,
+      float2* /* ptr */,
+      const float /* weight */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void
-  index_weighted_store(uint32_t idx, uint8_t* ptr, const float weight) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_weighted_store(
+      uint32_t /* idx */,
+      uint8_t* /* ptr */,
+      const float /* weight */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
 #if USE_ROCM_OR_CUDA_SM80_PLUS
-  DEVICE_INLINE void index_store(uint32_t idx, bfloat16_4* ptr) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_store(uint32_t /* idx */, bfloat16_4* /* ptr */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void
-  index_weighted_store(uint32_t idx, bfloat16_4* ptr, const float weight) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_weighted_store(
+      uint32_t /* idx */,
+      bfloat16_4* /* ptr */,
+      const float /* weight */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 #endif
 };
@@ -455,74 +471,86 @@ struct Vec4StepT<STEP, uint8_t> : Vec4AccT {
 template <uint32_t STEP>
 struct Vec4StepT<STEP, c10::Float8_e4m3fn> : Vec4AccT {
   DEVICE_INLINE Vec4StepT() {
-    CUDA_KERNEL_ASSERT(false);
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void load(const c10::Float8_e4m3fn* ptr, const uint32_t idx) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void load(
+      const c10::Float8_e4m3fn* /* ptr */,
+      const uint32_t /* idx */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
   DEVICE_INLINE void sum() {
-    CUDA_KERNEL_ASSERT(false);
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
   DEVICE_INLINE void weighted_sum(
-      const float* const weights,
-      const uint32_t idx_shift,
-      const uint32_t idx_scale) {
-    CUDA_KERNEL_ASSERT(false);
+      const float* const /* weights */,
+      const uint32_t /* idx_shift */,
+      const uint32_t /* idx_scale */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void index_add(uint32_t idx) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_add(uint32_t /* idx */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void index_fma(uint32_t idx, const float weight) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_fma(uint32_t /* idx */, const float /* weight */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void index_store(uint32_t idx, float4* ptr) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_store(uint32_t /* idx */, float4* /* ptr */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void index_store(uint32_t idx, float2* ptr) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_store(uint32_t /* idx */, float2* /* ptr */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void index_store(uint32_t idx, c10::Float8_e4m3fn* ptr) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_store(
+      uint32_t /* idx */,
+      c10::Float8_e4m3fn* /* ptr */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void
-  index_weighted_store(uint32_t idx, float4* ptr, const float weight) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_weighted_store(
+      uint32_t /* idx */,
+      float4* /* ptr */,
+      const float /* weight */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void
-  index_weighted_store(uint32_t idx, float2* ptr, const float weight) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_weighted_store(
+      uint32_t /* idx */,
+      float2* /* ptr */,
+      const float /* weight */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void
-  index_weighted_store(uint32_t idx, uint8_t* ptr, const float weight) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_weighted_store(
+      uint32_t /* idx */,
+      uint8_t* /* ptr */,
+      const float /* weight */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
   DEVICE_INLINE void index_weighted_store(
-      uint32_t idx,
-      c10::Float8_e4m3fn* ptr,
-      const float weight) {
-    CUDA_KERNEL_ASSERT(false);
+      uint32_t /* idx */,
+      c10::Float8_e4m3fn* /* ptr */,
+      const float /* weight */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
 #if USE_ROCM_OR_CUDA_SM80_PLUS
-  DEVICE_INLINE void index_store(uint32_t idx, bfloat16_4* ptr) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_store(uint32_t /* idx */, bfloat16_4* /* ptr */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void
-  index_weighted_store(uint32_t idx, bfloat16_4* ptr, const float weight) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_weighted_store(
+      uint32_t /* idx */,
+      bfloat16_4* /* ptr */,
+      const float /* weight */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 #endif
 };
@@ -530,74 +558,86 @@ struct Vec4StepT<STEP, c10::Float8_e4m3fn> : Vec4AccT {
 template <uint32_t STEP>
 struct Vec4StepT<STEP, c10::Float8_e4m3fnuz> : Vec4AccT {
   DEVICE_INLINE Vec4StepT() {
-    CUDA_KERNEL_ASSERT(false);
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void load(const c10::Float8_e4m3fnuz* ptr, const uint32_t idx) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void load(
+      const c10::Float8_e4m3fnuz* /* ptr */,
+      const uint32_t /* idx */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
   DEVICE_INLINE void sum() {
-    CUDA_KERNEL_ASSERT(false);
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
   DEVICE_INLINE void weighted_sum(
-      const float* const weights,
-      const uint32_t idx_shift,
-      const uint32_t idx_scale) {
-    CUDA_KERNEL_ASSERT(false);
+      const float* const /* weights */,
+      const uint32_t /* idx_shift */,
+      const uint32_t /* idx_scale */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void index_add(uint32_t idx) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_add(uint32_t /* idx */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void index_fma(uint32_t idx, const float weight) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_fma(uint32_t /* idx */, const float /* weight */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void index_store(uint32_t idx, float4* ptr) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_store(uint32_t /* idx */, float4* /* ptr */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void index_store(uint32_t idx, float2* ptr) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_store(uint32_t /* idx */, float2* /* ptr */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void index_store(uint32_t idx, c10::Float8_e4m3fnuz* ptr) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_store(
+      uint32_t /* idx */,
+      c10::Float8_e4m3fnuz* /* ptr */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void
-  index_weighted_store(uint32_t idx, float4* ptr, const float weight) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_weighted_store(
+      uint32_t /* idx */,
+      float4* /* ptr */,
+      const float /* weight */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void
-  index_weighted_store(uint32_t idx, float2* ptr, const float weight) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_weighted_store(
+      uint32_t /* idx */,
+      float2* /* ptr */,
+      const float /* weight */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void
-  index_weighted_store(uint32_t idx, uint8_t* ptr, const float weight) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_weighted_store(
+      uint32_t /* idx */,
+      uint8_t* /* ptr */,
+      const float /* weight */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
   DEVICE_INLINE void index_weighted_store(
-      uint32_t idx,
-      c10::Float8_e4m3fnuz* ptr,
-      const float weight) {
-    CUDA_KERNEL_ASSERT(false);
+      uint32_t /* idx */,
+      c10::Float8_e4m3fnuz* /* ptr */,
+      const float /* weight */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
 #if USE_ROCM_OR_CUDA_SM80_PLUS
-  DEVICE_INLINE void index_store(uint32_t idx, bfloat16_4* ptr) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_store(uint32_t /* idx */, bfloat16_4* /* ptr */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 
-  DEVICE_INLINE void
-  index_weighted_store(uint32_t idx, bfloat16_4* ptr, const float weight) {
-    CUDA_KERNEL_ASSERT(false);
+  DEVICE_INLINE void index_weighted_store(
+      uint32_t /* idx */,
+      bfloat16_4* /* ptr */,
+      const float /* weight */) {
+    CUDA_KERNEL_ASSERT(false && "Unsupported Vec4 operation");
   }
 #endif
 };
diff --git a/fbgemm_gpu/include/fbgemm_gpu/utils/vec_quant.cuh b/fbgemm_gpu/include/fbgemm_gpu/utils/vec_quant.cuh
index 0a25534e25..3d2336a45d 100644
--- a/fbgemm_gpu/include/fbgemm_gpu/utils/vec_quant.cuh
+++ b/fbgemm_gpu/include/fbgemm_gpu/utils/vec_quant.cuh
@@ -201,7 +201,7 @@ DEVICE_INLINE bfx4 fx4_to_bfx4(fx4 a) {
 
 template <typename T>
 DEVICE_INLINE T shfl_xor(
-    unsigned shfl_sync_mask,
+    [[maybe_unused]] unsigned shfl_sync_mask,
     const T val,
     int laneMask,
     int width = kThreadsPerWarp) {
diff --git a/fbgemm_gpu/src/split_embeddings_cache/lru_cache_populate_byte.cu b/fbgemm_gpu/src/split_embeddings_cache/lru_cache_populate_byte.cu
index e0ab983b5f..438a5f0090 100644
--- a/fbgemm_gpu/src/split_embeddings_cache/lru_cache_populate_byte.cu
+++ b/fbgemm_gpu/src/split_embeddings_cache/lru_cache_populate_byte.cu
@@ -271,7 +271,7 @@ __launch_bounds__(kMaxThreads) void direct_mapped_lru_cache_insert_byte_kernel(
     const pta::PackedTensorAccessor32<index_t, 1, at::RestrictPtrTraits>
         linear_cache_indices,
     pta::PackedTensorAccessor32<int64_t, 2, at::RestrictPtrTraits>
-        lxu_cache_miss_timestamp,
+        lxu_cache_miss_timestamp [[maybe_unused]],
     pta::PackedTensorAccessor32<int32_t, 1, at::RestrictPtrTraits> cache_sets,
     const bool gather_cache_stats,
     pta::PackedTensorAccessor32<int32_t, 1, at::RestrictPtrTraits>
diff --git a/fbgemm_gpu/src/split_embeddings_utils/transpose_embedding_input.cu b/fbgemm_gpu/src/split_embeddings_utils/transpose_embedding_input.cu
index 12819f97eb..429a2c8032 100644
--- a/fbgemm_gpu/src/split_embeddings_utils/transpose_embedding_input.cu
+++ b/fbgemm_gpu/src/split_embeddings_utils/transpose_embedding_input.cu
@@ -77,8 +77,8 @@ __global__ __launch_bounds__(kMaxThreads) void linearize_index_kernel(
         linear_indices,
     const int32_t info_B_num_bits,
     const uint32_t info_B_mask,
-    const uint32_t max_T,
-    const uint32_t max_B,
+    [[maybe_unused]] const uint32_t max_T,
+    [[maybe_unused]] const uint32_t max_B,
     // Use a raw pointer to avoid creating dummy PackedTensorAccessor
     const uint32_t* const __restrict__ vbe_b_t_map,
     FixedDivisor fd) {
diff --git a/include/fbgemm/ConvUtils.h b/include/fbgemm/ConvUtils.h
index 3b72ea52cf..c525e4d03b 100644
--- a/include/fbgemm/ConvUtils.h
+++ b/include/fbgemm/ConvUtils.h
@@ -61,7 +61,7 @@ struct conv_param_t {
       std::array<int, SPATIAL_DIM * 2> pd,
       std::array<int, SPATIAL_DIM> dilations = array_of_ones<SPATIAL_DIM>(),
       std::array<int, SPATIAL_DIM> otpt_pd = {},
-      bool transposed = false)
+      bool is_transposed = false)
       : MB(mb),
         IC(ic),
         OC(oc),
@@ -72,7 +72,7 @@ struct conv_param_t {
         pad(pd),
         dilation(dilations),
         output_pad(otpt_pd),
-        transposed(transposed) {
+        transposed(is_transposed) {
     if (ic % g != 0) {
       throw std::runtime_error(
           "groups = " + std::to_string(g) +
@@ -85,7 +85,7 @@ struct conv_param_t {
     }
 
     for (int d = 0; d < SPATIAL_DIM; ++d) {
-      if (transposed) {
+      if (is_transposed) {
         this->IN_DIMP[d] = this->IN_DIM[d] +
             (this->dilation[d] * (this->K[d] - 1) - this->pad[d]) +
             (this->dilation[d] * (this->K[d] - 1) - this->pad[SPATIAL_DIM + d]);
diff --git a/src/EmbeddingStatsTracker.h b/src/EmbeddingStatsTracker.h
index 8ae4cabe76..a8338ecd32 100644
--- a/src/EmbeddingStatsTracker.h
+++ b/src/EmbeddingStatsTracker.h
@@ -107,8 +107,8 @@ class EmbeddingStatsTracker {
      *
      * @param r Number of rows in the embedding table
      * @param d Embedding dimension
-     * @param batch_size Number of embeddings being looked up in a batch
-     * @param bag_size Number of embeddings pooled together (pooling factor)
+     * @param batch_sz Number of embeddings being looked up in a batch
+     * @param bag_sz Number of embeddings pooled together (pooling factor)
      * @param dt Data type used for storing the embedding values
      *
      * This constructor creates a unique entry that identifies an embedding
@@ -117,14 +117,14 @@ class EmbeddingStatsTracker {
     AccessPatternEntry(
         int64_t r,
         int64_t d,
-        int64_t batch_size,
-        int64_t bag_size,
+        int64_t batch_sz,
+        int64_t bag_sz,
         DataType input_dt,
         DataType output_dt)
         : rows(r),
           dims(d),
-          batch_size{batch_size},
-          bag_size(bag_size),
+          batch_size{batch_sz},
+          bag_size(bag_sz),
           input_data_type(input_dt),
           output_data_type(output_dt) {}