[GPU] INPUT/OUTPUT_LENGTH support shape agnostic kernel (#33434)

-  code clean-up
 - INPUT/OUTPUT_LENGTH in JIT is 0 at the shape agnostic kernel.

### Tickets:
     - N/A

Author: sungeunk

src/plugins/intel_gpu/src/kernel_selector/cl_kernels/random_uniform_ref.cl

@@ -132,7 +132,7 @@ KERNEL(random_uniform_ref)(OPTIONAL_SHAPE_INFO_ARG
                             const __global INPUT2_TYPE* max_val, __global OUTPUT_TYPE *output) {
     const uint plain_index = get_global_id(0);
     uint4 result = FUNC_CALL(run_philox)(plain_index);
-    FILL_FUNC(FUNC_NAME(OUTPUT_TYPE), result, *min_val, *max_val, output, plain_index * OUTPUT_STEP, COMPUTATIONAL_OPERATIONS_NUMBER);
+    FILL_FUNC(FUNC_NAME(OUTPUT_TYPE), result, *min_val, *max_val, output, plain_index * OUTPUT_STEP, OUTPUT_LENGTH);
 }
 
 #undef FILL_FUNC

src/plugins/intel_gpu/src/kernel_selector/cl_kernels/reduce_gpu_b_fs_yx_fsv16.cl

@@ -190,7 +190,7 @@ KERNEL(reduce_fsv16)(
 #endif
 
     const uint linear_idx = FUNC_CALL(calc_linear_offset)(b, f, y, x);
-    if (linear_idx >= COMPUTATIONAL_OPERATIONS_NUMBER)
+    if (linear_idx >= OUTPUT_LENGTH)
         return;
 
     const uint input_x_pitch = FSV;

src/plugins/intel_gpu/src/kernel_selector/cl_kernels/reduce_ref.cl

@@ -75,7 +75,7 @@ KERNEL(reduce_ref)(
 #endif
 
     const uint linear_idx = FUNC_CALL(calc_linear_offset)(OPTIONAL_SHAPE_INFO_TENSOR b, f, v, u, w, z, y, x);
-    if (linear_idx >= COMPUTATIONAL_OPERATIONS_NUMBER)
+    if (linear_idx >= OUTPUT_LENGTH)
         return;
 
 #ifdef REDUCE_BATCH

src/plugins/intel_gpu/src/kernel_selector/jitter.cpp

@@ -259,7 +259,6 @@ class TensorBaseTJitConstant : public JitConstant {
     JitDefinitions GetDefinitions(const Tensor::TensorBaseT<DType, Layout>& t) const {
         JitDefinitions definitions{
             {_name + "_VIEW_OFFSET", toCodeString(t.GetViewOffset())},
-            {_name + "_LENGTH", toCodeString(t.LogicalSize())},
             {_name + "_DIMS", toCodeString(t.GetDims().size())},
             {_name + "_SIMPLE", toCodeString(t.SimpleLayout())},
             {_name + "_GROUPED", toCodeString(t.GroupedLayout())},
@@ -363,6 +362,9 @@ JitDefinitions DataTensorJitConstant::GetDefinitions() const {
         {_name + "_PAD_AFTER_BATCH_NUM", dims_padded.b_pad().second},
     };
     if (_tensor.is_dynamic()) {
+        const std::string total_data_size = toVectorMulString({dims.x(), dims.y(), dims.z(), dims.w(), dims.f(), dims.b()});
+        definitions.push_back({_name + "_LENGTH", toCodeString(total_data_size)});
+
         if (_tensor.GetLayout() == DataLayout::bf || _tensor.GetLayout() == DataLayout::bfyx ||
             _tensor.GetLayout() == DataLayout::bfzyx || _tensor.GetLayout() == DataLayout::bfwzyx ||
             _tensor.GetLayout() == DataLayout::bfuwzyx || _tensor.GetLayout() == DataLayout::bfvuwzyx ||
@@ -394,6 +396,8 @@ JitDefinitions DataTensorJitConstant::GetDefinitions() const {
             OPENVINO_ASSERT(false, "[GPU] Jitter couldn't generate dynamic pitches for given layout");
         }
     } else {
+        definitions.push_back({_name + "_LENGTH", toCodeString(_tensor.LogicalSize())});
+
         // static dim
         definitions.push_back({_name + "_X_PITCH", toCodeString(_tensor.X().pitch)});
         definitions.push_back({_name + "_Y_PITCH", toCodeString(_tensor.Y().pitch)});
@@ -938,6 +942,7 @@ JitDefinitions WeightTensorJitConstant::GetDefinitions() const {
         {_name + "_IFM_PITCH", toCodeString(_tensor.IFM().pitch)},
         {_name + "_OFM_PITCH", toCodeString(_tensor.OFM().pitch)},
         {_name + "_GROUPS_PITCH", toCodeString(_tensor.G().pitch)},
+        {_name + "_LENGTH", toCodeString(_tensor.LogicalSize())},
     };
 
     definitions.insert(definitions.end(), baseDefinitions.begin(), baseDefinitions.end());

src/plugins/intel_gpu/src/kernel_selector/kernels/random_uniform/random_uniform_kernel_ref.cpp

@@ -40,21 +40,6 @@ JitConstants RandomUniformKernelRef::GetJitConstants(const random_uniform_params
         jit.AddConstant(MakeJitConstant("GLOBAL_SEED", params.global_seed));
     }
 
-    const auto& output = params.outputs[0];
-    if (output.is_dynamic()) {
-        DimensionAccessHelperJit dims(output);
-        jit.AddConstant(MakeJitConstant("COMPUTATIONAL_OPERATIONS_NUMBER", toVectorMulString({dims.x(),
-                                                                                              dims.y(),
-                                                                                              dims.z(),
-                                                                                              dims.w(),
-                                                                                              dims.u(),
-                                                                                              dims.v(),
-                                                                                              dims.f(),
-                                                                                              dims.b()})));
-    } else {
-        jit.AddConstant(MakeJitConstant("COMPUTATIONAL_OPERATIONS_NUMBER", params.outputs[0].LogicalSize()));
-    }
-
     jit.AddConstant(MakeJitConstant("OP_SEED", params.op_seed));
     jit.AddConstant(MakeJitConstant("OUTPUT_STEP", getStep(params)));
     return jit;

src/plugins/intel_gpu/src/kernel_selector/kernels/reduce/reduce_kernel_base.cpp

@@ -27,22 +27,6 @@ bool ReduceKernelBase::Validate(const Params& p) const {
 
 JitConstants ReduceKernelBase::GetJitConstants(const reduce_params& params) const {
     JitConstants jit = MakeBaseParamsJitConstants(params);
-
-    const auto& output = params.outputs[0];
-    if (output.is_dynamic()) {
-        DimensionAccessHelperJit dims(output);
-        jit.AddConstant(MakeJitConstant("COMPUTATIONAL_OPERATIONS_NUMBER", toVectorMulString({dims.x(),
-                                                                                              dims.y(),
-                                                                                              dims.z(),
-                                                                                              dims.w(),
-                                                                                              dims.u(),
-                                                                                              dims.v(),
-                                                                                              dims.f(),
-                                                                                              dims.b()})));
-    } else {
-        jit.AddConstant(MakeJitConstant("COMPUTATIONAL_OPERATIONS_NUMBER", params.outputs[0].LogicalSize()));
-    }
-
     jit.AddConstant(MakeJitConstant("REDUCE_" + toString(params.reduceMode) + "_MODE", 1));
     jit.AddConstant(MakeJitConstant("KEEP_DIMS", params.keepDims));