[Common] Fix long compile time in padding.cu on arch 75 (#2562)

* Fix long compile time in padding.cu

Signed-off-by: Jeremy Berchtold <[email protected]>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

---------

Signed-off-by: Jeremy Berchtold <[email protected]>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>

Author: jberchtold-nvidia

transformer_engine/common/util/padding.cu

@@ -94,15 +94,14 @@ __global__ void __launch_bounds__(threads_per_block) multi_padding_kernel(MultiP
 #pragma unroll
     for (int i2 = 0; i2 < nvec; ++i2) {
       const int row = tile_row + i1 * nvec + i2;
-      size_t row_offset = static_cast<size_t>(row) * row_length;
       const int col = tile_col + j1 * nvec;
       Vec local_input;
       Vec local_output;
       local_input.clear();
       if (row < num_rows) {
         for (int j2 = 0; j2 < nvec; ++j2) {
           if (col + j2 < row_length) {
-            local_input.data.elt[j2] = input[row_offset + col + j2];
+            local_input.data.elt[j2] = input[static_cast<size_t>(row) * row_length + col + j2];
           }
         }
       }
@@ -113,14 +112,14 @@ __global__ void __launch_bounds__(threads_per_block) multi_padding_kernel(MultiP
       if (row < num_rows) {
         for (int j2 = 0; j2 < nvec; ++j2) {
           if (col + j2 < row_length) {
-            output[row_offset + col + j2] = local_output.data.elt[j2];
+            output[static_cast<size_t>(row) * row_length + col + j2] = local_output.data.elt[j2];
           }
         }
       } else if (row < padded_num_rows) {
         // padding
         for (int j2 = 0; j2 < nvec; ++j2) {
           if (col + j2 < row_length) {
-            output[row_offset + col + j2] = local_zero;
+            output[static_cast<size_t>(row) * row_length + col + j2] = local_zero;
           }
         }
       }
@@ -179,15 +178,14 @@ __global__ void __launch_bounds__(threads_per_block) multi_unpadding_kernel(Mult
 #pragma unroll
     for (int i2 = 0; i2 < nvec; ++i2) {
       const int row = tile_row + i1 * nvec + i2;
-      size_t row_offset = static_cast<size_t>(row) * row_length;
       const int col = tile_col + j1 * nvec;
       Vec local_input;
       Vec local_output;
       local_input.clear();
       if (row < num_rows) {
         for (int j2 = 0; j2 < nvec; ++j2) {
           if (col + j2 < row_length) {
-            local_input.data.elt[j2] = input[row_offset + col + j2];
+            local_input.data.elt[j2] = input[static_cast<size_t>(row) * row_length + col + j2];
           }
         }
       }
@@ -198,7 +196,7 @@ __global__ void __launch_bounds__(threads_per_block) multi_unpadding_kernel(Mult
       if (row < num_rows) {
         for (int j2 = 0; j2 < nvec; ++j2) {
           if (col + j2 < row_length) {
-            output[row_offset + col + j2] = local_output.data.elt[j2];
+            output[static_cast<size_t>(row) * row_length + col + j2] = local_output.data.elt[j2];
           }
         }
       }