use a concise expression to further speedup the cuda kernel

Author: bugparty

ggml/src/ggml-cuda/pad_reflect_1d.cu

@@ -36,20 +36,20 @@ static __global__ __launch_bounds__(CUDA_PAD_REFLECT_1D_BLOCK_SIZE, 1) void
     const char * src0_ptr = (const char *) src0 + i3 * nb03 + i2 * nb02 + i1 * nb01;
     char *       dst_ptr  = (char *) dst + i3 * nb3 + i2 * nb2 + i1 * nb1;
 
-    float         value;
-    const int64_t j = i0 - p0;
+    const int64_t rel_i0 = i0 - p0;
+    int64_t src_idx;
 
-    if (j < 0) {  // i0<p0
+    if (rel_i0 < 0) {
         // Left padding - reflect
-        value = *(const float *) (src0_ptr - j * nb00);
-    } else if (j < ne00) {  //i0 < ne0 - p1
+        src_idx = -rel_i0;
+    } else if (rel_i0 < ne00) {
         // Middle - copy
-        value = *(const float *) (src0_ptr + j * nb00);
+        src_idx = rel_i0;
     } else {
         // Right padding - reflect
-        const int64_t src_idx = (ne0 - p1 - p0) - (p1 + 1 - (ne0 - i0)) - 1;
-        value                 = *(const float *) (src0_ptr + src_idx * nb00);
+        src_idx = 2 * ne00 - 2 - rel_i0;
     }
+    const float value               = *(const float *) (src0_ptr + src_idx * nb00);
     *(float *) (dst_ptr + i0 * nb0) = value;
 }
 

tests/test-backend-ops.cpp

@@ -6636,8 +6636,8 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
 
     test_cases.emplace_back(new test_pad_reflect_1d(GGML_TYPE_F32, {512, 34, 2, 1}));
     test_cases.emplace_back(new test_pad_reflect_1d(GGML_TYPE_F32, {3000, 80, 1, 1}));
-    test_cases.emplace_back(new test_pad_reflect_1d(GGML_TYPE_F32, {3000, 384, 1, 1}));
     test_cases.emplace_back(new test_pad_reflect_1d(GGML_TYPE_F32, {3000, 80, 4, 1}));
+    test_cases.emplace_back(new test_pad_reflect_1d(GGML_TYPE_F32, {3000, 384, 1, 1}));
     test_cases.emplace_back(new test_pad_reflect_1d(GGML_TYPE_F32, {3000, 384, 4, 1}));
 
     test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F16, GGML_TYPE_F32, 16416, 1, 128, {8,  1}, {4, 1}, {0, 2, 1, 3}));