Promote -Wzero-as-null-pointer-constant (#6336)

cmake/modules/CppLibrary.cmake

@@ -179,15 +179,7 @@ function(fbgemm_get_warning_flags)
     -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-error=attributes)
 
   # 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

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

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

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());

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);

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");
 }
 

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);

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);

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.

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;
 }