eliminate warning of process_block

liligwu · Bernard-Liu · commit 813265cd738e · 2025-11-13T17:35:22.000Z
diff --git a/fbgemm_gpu/codegen/training/backward/embedding_backward_split_device_kernel_template.cuh b/fbgemm_gpu/codegen/training/backward/embedding_backward_split_device_kernel_template.cuh
@@ -237,17 +237,13 @@ DEVICE_INLINE void compute_grad_sum_{{ kdesc }}(
 
             // Process blocks of different sizes with loop unrolling
             if constexpr (sizeof(grad_t) <= 2) {
-                #pragma unroll kFixedMaxVecsPerThread
                 PROCESS_BLOCK(8, kFixedMaxVecsPerThread, grad_sum, grad_output, grad_offset, \
                     vec_start, kThreadGroupSize, threadIdx.x, VEC_WIDTH, D, j, sl, sl_end)
             }
-            #pragma unroll kFixedMaxVecsPerThread
             PROCESS_BLOCK(4, kFixedMaxVecsPerThread, grad_sum, grad_output, grad_offset, \
                 vec_start, kThreadGroupSize, threadIdx.x, VEC_WIDTH, D, j, sl, sl_end)
-            #pragma unroll kFixedMaxVecsPerThread
             PROCESS_BLOCK(2, kFixedMaxVecsPerThread, grad_sum, grad_output, grad_offset, \
                 vec_start, kThreadGroupSize, threadIdx.x, VEC_WIDTH, D, j, sl, sl_end)
-            #pragma unroll kFixedMaxVecsPerThread
             PROCESS_BLOCK(1, kFixedMaxVecsPerThread, grad_sum, grad_output, grad_offset, \
                 vec_start, kThreadGroupSize, threadIdx.x, VEC_WIDTH, D, j, sl, sl_end)
         }
@@ -266,6 +262,7 @@ DEVICE_INLINE void compute_grad_sum_{{ kdesc }}(
     }
 }
 
+#undef PROCESS_BLOCK
 {%- endif %}
 
     // clang-format on

Original file line number	Diff line number	Diff line change
`@@ -237,17 +237,13 @@ DEVICE_INLINE void compute_grad_sum_{{ kdesc }}(`
`237`	`237`
`238`	`238`	`// Process blocks of different sizes with loop unrolling`
`239`	`239`	`if constexpr (sizeof(grad_t) <= 2) {`
`240`		`- #pragma unroll kFixedMaxVecsPerThread`
`241`	`240`	`PROCESS_BLOCK(8, kFixedMaxVecsPerThread, grad_sum, grad_output, grad_offset, \`
`242`	`241`	`vec_start, kThreadGroupSize, threadIdx.x, VEC_WIDTH, D, j, sl, sl_end)`
`243`	`242`	`}`
`244`		`- #pragma unroll kFixedMaxVecsPerThread`
`245`	`243`	`PROCESS_BLOCK(4, kFixedMaxVecsPerThread, grad_sum, grad_output, grad_offset, \`
`246`	`244`	`vec_start, kThreadGroupSize, threadIdx.x, VEC_WIDTH, D, j, sl, sl_end)`
`247`		`- #pragma unroll kFixedMaxVecsPerThread`
`248`	`245`	`PROCESS_BLOCK(2, kFixedMaxVecsPerThread, grad_sum, grad_output, grad_offset, \`
`249`	`246`	`vec_start, kThreadGroupSize, threadIdx.x, VEC_WIDTH, D, j, sl, sl_end)`
`250`		`- #pragma unroll kFixedMaxVecsPerThread`
`251`	`247`	`PROCESS_BLOCK(1, kFixedMaxVecsPerThread, grad_sum, grad_output, grad_offset, \`
`252`	`248`	`vec_start, kThreadGroupSize, threadIdx.x, VEC_WIDTH, D, j, sl, sl_end)`
`253`	`249`	`}`
`@@ -266,6 +262,7 @@ DEVICE_INLINE void compute_grad_sum_{{ kdesc }}(`
`266`	`262`	`}`
`267`	`263`	`}`
`268`	`264`
	`265`	`+#undef PROCESS_BLOCK`
`269`	`266`	`{%- endif %}`
`270`	`267`
`271`	`268`	`// clang-format on`