[GPU] Generate a new RandomUniform sequence on every inference

src/plugins/intel_gpu/src/graph/graph_optimizer/mark_shape_of_subgraphs.cpp

@@ -15,6 +15,7 @@
 #include "pass_manager.h"
 #include "fully_connected_inst.h"
 #include "gemm_inst.h"
+#include "random_uniform_inst.h"
 
 #include "intel_gpu/graph/program.hpp"
 
@@ -87,6 +88,11 @@ bool mark_shape_of_subgraphs::can_mark_node(const program_node& node) {
         return false;
     }
 
+    // RandomUniform output changes on every inference, while shape_of subgraphs are skipped when input shapes do not change
+    if (node.is_type<random_uniform>()) {
+        return false;
+    }
+
     // CPU implementation does not support float data types for mask and mixed types for data inputs, so check them
     // before including it into shape_of sub-graph
     if (node.is_type<select>() &&

src/plugins/intel_gpu/src/graph/impls/ocl/random_uniform.cpp

@@ -32,6 +32,24 @@ struct random_uniform_impl : typed_primitive_impl_ocl<random_uniform> {
         }
     }
 
+    event::ptr execute_impl(const std::vector<event::ptr>& events, random_uniform_inst& instance) override {
+        // RandomUniform must produce a new sequence on every execution: pass the current Philox state
+        // as runtime scalars, re-set kernel arguments (static graphs set them only once), then advance the state.
+        if (!instance.can_be_optimized()) {
+            OPENVINO_ASSERT(_kernel_data.kernels.size() == 1 && _kernel_data.kernels[0].params.scalars.size() == 3,
+                            "[GPU] RandomUniform kernel expects 3 runtime scalars (global_seed, n_state, counter_state)");
+            auto& scalars = _kernel_data.kernels[0].params.scalars;
+            scalars[0].v.u64 = instance.get_global_seed();
+            scalars[1].v.u64 = instance.get_n_state();
+            scalars[2].v.u64 = instance.get_counter_state();
+            parent::set_arguments_impl(instance);
+        }
+
+        auto ev = parent::execute_impl(events, instance);
+        instance.advance_state(instance.get_output_layout().count());
+        return ev;
+    }
+
 public:
     static kernel_params_t get_kernel_params(const kernel_impl_params& impl_param, bool is_shape_agnostic = false) {
         const auto& primitive = impl_param.typed_desc<random_uniform>();

src/plugins/intel_gpu/src/graph/include/random_uniform_inst.h

@@ -35,6 +35,22 @@ class typed_primitive_inst<random_uniform> : public typed_primitive_inst_base<ra
     static std::string to_string(random_uniform_node const &node);
 
     typed_primitive_inst(network &network, random_uniform_node const &desc);
+
+    // Philox state. It is advanced after every execution so that each inference produces a new sequence.
+    uint64_t get_global_seed() const {
+        return _global_seed;
+    }
+    uint64_t get_n_state() const {
+        return _n_state;
+    }
+    uint64_t get_counter_state() const;
+    void advance_state(size_t generated_elements);
+
+private:
+    uint64_t _global_seed = 0;
+    uint64_t _op_seed = 0;
+    uint64_t _n_state = 0;
+    uint64_t _counter_state = 0;
 };
 
 using random_uniform_inst = typed_primitive_inst<random_uniform>;

src/plugins/intel_gpu/src/graph/program.cpp

@@ -41,6 +41,7 @@
 #include "reorg_yolo_inst.h"
 #include "eltwise_inst.h"
 #include "non_zero_inst.h"
+#include "random_uniform_inst.h"
 #include "softmax_inst.h"
 #include "permute_inst.h"
 #include "custom_gpu_primitive_inst.h"
@@ -651,7 +652,9 @@ void program::post_optimize_graph(bool is_internal) {
 
 // mark if the node is constant assuming that all dependencies are marked properly
 void program::mark_if_constant(program_node& node) {
-    if (node.get_dependencies().empty() || node.is_type<assign>() || node.is_type<read_value>() || node.is_type<gather_nonzero>()) {
+    // RandomUniform must produce a new sequence on every inference, so it must never be folded into a constant
+    if (node.get_dependencies().empty() || node.is_type<assign>() || node.is_type<read_value>() || node.is_type<gather_nonzero>() ||
+        node.is_type<random_uniform>()) {
         return;
     }
     node.constant = true;

src/plugins/intel_gpu/src/graph/random_uniform.cpp

@@ -6,13 +6,35 @@
 #include "primitive_type_base.h"
 #include "json_object.h"
 
+#include <random>
+
 #include "random_uniform_shape_inference.hpp"
 
 namespace cldnn {
 GPU_DEFINE_PRIMITIVE_TYPE_ID(random_uniform)
 
 random_uniform_inst::typed_primitive_inst(network& network, random_uniform_node const &node)
 : parent(network, node) {
+    const auto& desc = node.get_primitive();
+    _global_seed = desc->global_seed;
+    _op_seed = desc->op_seed;
+    // When both seed values are equal to zero RandomUniform should generate non-deterministic sequence.
+    if (_global_seed == 0 && _op_seed == 0) {
+        _global_seed = std::random_device{}();
+    }
+}
+
+uint64_t random_uniform_inst::get_counter_state() const {
+    return _counter_state > 0 ? _counter_state : _op_seed;
+}
+
+void random_uniform_inst::advance_state(size_t generated_elements) {
+    // Same as the CPU plugin: skip 256 * <number of output elements> values of n, carry the overflow into the counter.
+    const uint64_t skip_count = static_cast<uint64_t>(generated_elements) * 256;
+    _n_state += skip_count;
+    if (_n_state < skip_count) {
+        _counter_state++;
+    }
 }
 
 layout random_uniform_inst::calc_output_layout(random_uniform_node const &node, kernel_impl_params const& impl_param) {

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

@@ -42,9 +42,7 @@ inline ulong FUNC(raise_key)(uint2 key_lr) {
     return FUNC_CALL(unite_high_low)(key_lr[1], key_lr[0]);
 }
 
-inline uint4 FUNC(run_philox)(ulong n) {
-    ulong key = GLOBAL_SEED;
-    ulong counter = OP_SEED;
+inline uint4 FUNC(run_philox)(ulong key, ulong counter, ulong n) {
     ulong2 counter_n = {counter, n};
     for (size_t i = 0; i < N_ROUNDS; i++) {
         counter_n = FUNC_CALL(calculate_round)(key, counter_n);
@@ -151,9 +149,14 @@ inline void FUNC(fill_bf16)(const uint4 res,
 
 KERNEL(random_uniform_ref)(OPTIONAL_SHAPE_INFO_ARG
                             const __global INPUT0_TYPE* shape, const __global INPUT1_TYPE *min_val,
-                            const __global INPUT2_TYPE* max_val, __global OUTPUT_TYPE *output) {
+                            const __global INPUT2_TYPE* max_val, __global OUTPUT_TYPE *output,
+                            const ulong global_seed, const ulong n_state, const ulong counter_state) {
     const uint plain_index = get_global_id(0);
-    uint4 result = FUNC_CALL(run_philox)(plain_index);
+    // Philox state is passed at runtime and advanced by the host after every execution,
+    // so that each inference produces a new sequence (same scheme as the CPU plugin).
+    const ulong n = n_state + plain_index;
+    const ulong counter = counter_state + (n < n_state ? 1 : 0);
+    uint4 result = FUNC_CALL(run_philox)(global_seed, counter, n);
     FILL_FUNC(FUNC_NAME(OUTPUT_TYPE_NAME), result, DECODE_INPUT1_COMPUTE_TYPE(*min_val), DECODE_INPUT2_COMPUTE_TYPE(*max_val), output, plain_index * OUTPUT_STEP, OUTPUT_LENGTH);
 }
 

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

@@ -5,7 +5,6 @@
 #include "random_uniform_kernel_ref.h"
 
 #include <kernel_selector_utils.h>
-#include <random>
 
 
 namespace kernel_selector {
@@ -33,14 +32,8 @@ CommonDispatchData RandomUniformKernelRef::SetDefault(const random_uniform_param
 
 JitConstants RandomUniformKernelRef::GetJitConstants(const random_uniform_params &params) const {
     JitConstants jit = MakeBaseParamsJitConstants(params);
-    // When both seed values are equal to zero RandomUniform should generate non-deterministic sequence.
-    if (params.global_seed == 0 && params.op_seed == 0) {
-        jit.AddConstant(MakeJitConstant("GLOBAL_SEED", std::random_device {}()));
-    } else {
-        jit.AddConstant(MakeJitConstant("GLOBAL_SEED", params.global_seed));
-    }
-
-    jit.AddConstant(MakeJitConstant("OP_SEED", params.op_seed));
+    // Seeds and the Philox state are runtime scalar arguments (see random_uniform_inst), not JIT constants:
+    // the state must advance between inferences, and the kernel binary must not depend on the seed.
     jit.AddConstant(MakeJitConstant("OUTPUT_STEP", getStep(params)));
 
     if (params.outputs[0].GetDType() == Datatype::BF16) {
@@ -71,6 +64,16 @@ KernelsData RandomUniformKernelRef::GetKernelsData(const Params &params) const {
                      false, static_cast<int>(new_params.inputs.size()), static_cast<uint32_t>(new_params.fused_ops.size()),
                      static_cast<int>(new_params.outputs.size()), new_params.is_shape_agnostic);
 
+    // global_seed, n_state, counter_state. Values are set before every execution by the primitive impl.
+    auto& kernel = kernel_data.kernels[0];
+    for (uint32_t i = 0; i < 3; ++i) {
+        kernel.params.arguments.push_back({ArgumentDescriptor::Types::SCALAR, i});
+        ScalarDescriptor s;
+        s.t = ScalarDescriptor::Types::UINT64;
+        s.v.u64 = 0;
+        kernel.params.scalars.push_back(s);
+    }
+
     KernelsData kernelsData;
     kernelsData.push_back(std::move(kernel_data));
     return kernelsData;

src/plugins/intel_gpu/tests/functional/sources.cmake

@@ -358,6 +358,7 @@ set(GPU_FUNCTIONAL_TEST_SOURCES
     ${CMAKE_CURRENT_LIST_DIR}/subgraph_tests/matmul_sum_fusion.cpp
     ${CMAKE_CURRENT_LIST_DIR}/subgraph_tests/mvn_layernorm_byxf.cpp
     ${CMAKE_CURRENT_LIST_DIR}/subgraph_tests/prevent_f32_dequantize.cpp
+    ${CMAKE_CURRENT_LIST_DIR}/subgraph_tests/random_uniform_infer.cpp
     ${CMAKE_CURRENT_LIST_DIR}/subgraph_tests/rms_dynamic_quantize_fusion.cpp
     ${CMAKE_CURRENT_LIST_DIR}/subgraph_tests/sdpa.cpp
     ${CMAKE_CURRENT_LIST_DIR}/subgraph_tests/serialize.cpp

src/plugins/intel_gpu/tests/functional/subgraph_tests/random_uniform_infer.cpp

@@ -0,0 +1,129 @@
+// Copyright (C) 2018-2026 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#include <gtest/gtest.h>
+
+#include <vector>
+
+#include "common_test_utils/test_constants.hpp"
+#include "openvino/op/add.hpp"
+#include "openvino/op/constant.hpp"
+#include "openvino/op/parameter.hpp"
+#include "openvino/op/random_uniform.hpp"
+#include "openvino/op/result.hpp"
+#include "openvino/op/shape_of.hpp"
+#include "openvino/reference/random_uniform.hpp"
+#include "openvino/runtime/core.hpp"
+
+namespace {
+// RandomUniform with constant inputs must not be folded into a constant by the GPU plugin, and every
+// inference must produce a new sequence (the Philox state advances like in the CPU plugin).
+using RandomUniformInferParams = std::tuple<uint64_t /* global_seed */, uint64_t /* op_seed */>;
+
+class RandomUniformInferGPUTest : public ::testing::TestWithParam<RandomUniformInferParams> {
+public:
+    static std::string getTestCaseName(const ::testing::TestParamInfo<RandomUniformInferParams>& obj) {
+        return "global_seed=" + std::to_string(std::get<0>(obj.param)) + "_op_seed=" + std::to_string(std::get<1>(obj.param));
+    }
+
+protected:
+    const ov::Shape out_shape{1, 50, 32};
+
+    std::shared_ptr<ov::Model> make_model(uint64_t global_seed, uint64_t op_seed) const {
+        auto x = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, out_shape);
+        auto shape = ov::op::v0::Constant::create(ov::element::i64, {out_shape.size()}, out_shape);
+        auto min_val = ov::op::v0::Constant::create(ov::element::f32, {}, {0.0f});
+        auto max_val = ov::op::v0::Constant::create(ov::element::f32, {}, {1.0f});
+        auto rnd = std::make_shared<ov::op::v8::RandomUniform>(shape, min_val, max_val, ov::element::f32, global_seed, op_seed);
+        auto add = std::make_shared<ov::op::v1::Add>(x, rnd);
+        return std::make_shared<ov::Model>(ov::OutputVector{add}, ov::ParameterVector{x});
+    }
+
+    static std::vector<float> infer(ov::InferRequest& req, const ov::Tensor& input) {
+        req.set_input_tensor(input);
+        req.infer();
+        auto out = req.get_output_tensor();
+        return std::vector<float>(out.data<float>(), out.data<float>() + out.get_size());
+    }
+};
+
+TEST_P(RandomUniformInferGPUTest, NewSequenceOnEveryInference) {
+    const auto [global_seed, op_seed] = GetParam();
+    ov::Core core;
+    auto compiled = core.compile_model(make_model(global_seed, op_seed), ov::test::utils::DEVICE_GPU);
+
+    ov::Tensor input(ov::element::f32, out_shape);
+    std::fill_n(input.data<float>(), input.get_size(), 0.0f);
+
+    const size_t iterations = 3;
+    auto req = compiled.create_infer_request();
+    std::vector<std::vector<float>> results;
+    for (size_t i = 0; i < iterations; ++i) {
+        results.push_back(infer(req, input));
+    }
+    for (size_t i = 1; i < iterations; ++i) {
+        ASSERT_NE(results[i - 1], results[i]) << "RandomUniform returned the same values on inference " << i;
+    }
+
+    if (global_seed != 0 || op_seed != 0) {
+        // Deterministic seeds: each inference must match the sequence produced by the CPU plugin.
+        std::vector<uint64_t> out_shape_u64(out_shape.begin(), out_shape.end());
+        const float min_v = 0.0f, max_v = 1.0f;
+        std::pair<uint64_t, uint64_t> state{0, 0};
+        for (size_t it = 0; it < iterations; ++it) {
+            std::vector<float> expected(ov::shape_size(out_shape));
+            ov::reference::random_uniform(out_shape_u64.data(),
+                                          reinterpret_cast<const char*>(&min_v),
+                                          reinterpret_cast<const char*>(&max_v),
+                                          reinterpret_cast<char*>(expected.data()),
+                                          ov::Shape{out_shape.size()},
+                                          ov::element::f32,
+                                          global_seed,
+                                          op_seed,
+                                          state);
+            // Same state progression as the CPU plugin: n += 256 * <number of output elements>.
+            const uint64_t skip_count = ov::shape_size(out_shape) * 256;
+            state.first += skip_count;
+            if (state.first < skip_count) {
+                state.second++;
+            }
+            for (size_t i = 0; i < expected.size(); ++i) {
+                ASSERT_NEAR(expected[i], results[it][i], 1e-6f) << "inference " << it << ", i = " << i;
+            }
+        }
+    }
+}
+
+INSTANTIATE_TEST_SUITE_P(smoke_RandomUniformInfer,
+                         RandomUniformInferGPUTest,
+                         ::testing::Values(RandomUniformInferParams{0, 0}, RandomUniformInferParams{80, 100}),
+                         RandomUniformInferGPUTest::getTestCaseName);
+
+// Dynamic input with RandomUniform shape taken from ShapeOf(input), as produced by torch.randn_like conversion.
+// RandomUniform must not be skipped as a part of a shape_of subgraph when the input shape does not change.
+TEST(smoke_RandomUniformDynamicInferGPUTest, ShapeOfInputNewSequenceOnEveryInference) {
+    auto x = std::make_shared<ov::op::v0::Parameter>(ov::element::f32, ov::PartialShape::dynamic(3));
+    auto shape = std::make_shared<ov::op::v3::ShapeOf>(x);
+    auto min_val = ov::op::v0::Constant::create(ov::element::f32, {}, {0.0f});
+    auto max_val = ov::op::v0::Constant::create(ov::element::f32, {}, {1.0f});
+    auto rnd = std::make_shared<ov::op::v8::RandomUniform>(shape, min_val, max_val, ov::element::f32, 0, 0);
+    auto add = std::make_shared<ov::op::v1::Add>(x, rnd);
+    auto model = std::make_shared<ov::Model>(ov::OutputVector{add}, ov::ParameterVector{x});
+
+    ov::Core core;
+    auto req = core.compile_model(model, ov::test::utils::DEVICE_GPU).create_infer_request();
+    ov::Tensor input(ov::element::f32, ov::Shape{1, 50, 32});
+    std::fill_n(input.data<float>(), input.get_size(), 0.0f);
+
+    std::vector<std::vector<float>> results;
+    for (size_t i = 0; i < 3; ++i) {
+        req.set_input_tensor(input);
+        req.infer();
+        auto out = req.get_output_tensor();
+        results.emplace_back(out.data<float>(), out.data<float>() + out.get_size());
+    }
+    ASSERT_NE(results[0], results[1]);
+    ASSERT_NE(results[1], results[2]);
+}
+}  // namespace

src/plugins/intel_gpu/tests/unit/passes/mark_shape_of_subgraphs_test.cpp

@@ -21,6 +21,7 @@
 #include "paged_attention_inst.h"
 #include "fully_connected_inst.h"
 #include "gemm_inst.h"
+#include "random_uniform_inst.h"
 #include "pass_manager.h"
 #include "to_string_utils.h"
 
@@ -658,3 +659,47 @@ TEST(mark_shape_of_subgraphs, gemm_not_marked) {
     ASSERT_FALSE(prog->get_node("relu").is_in_shape_of_subgraph());
 }
 
+
+// RandomUniform produces a new sequence on every inference, so it must not be a part of a shape_of subgraph:
+// such subgraphs are skipped when the input shapes do not change, which would freeze the random values.
+TEST(mark_shape_of_subgraphs, random_uniform_no_mark) {
+    auto& engine = get_test_engine();
+    auto input_layout_dynamic = layout{ov::PartialShape{ov::Dimension::dynamic(), ov::Dimension::dynamic()},
+                                       data_types::f32, format::bfyx};
+    auto min_val = engine.allocate_memory({ ov::PartialShape{1}, data_types::f32, format::bfyx });
+    auto max_val = engine.allocate_memory({ ov::PartialShape{1}, data_types::f32, format::bfyx });
+    set_values(min_val, {0.f});
+    set_values(max_val, {1.f});
+
+    topology topology;
+    topology.add(input_layout("input", input_layout_dynamic));
+    topology.add(data("min_val", min_val));
+    topology.add(data("max_val", max_val));
+    topology.add(shape_of("shape_of", input_info("input"), data_types::i32));
+    topology.add(random_uniform("random_uniform", {input_info("shape_of"), input_info("min_val"), input_info("max_val")},
+                                data_types::f32, 0, 0));
+    topology.add(eltwise("eltwise", input_info("input"), input_info("random_uniform"), eltwise_mode::sum));
+
+    ExecutionConfig config = get_test_default_config(engine);
+    config.set_property(ov::intel_gpu::allow_new_shape_infer(true));
+    config.set_property(ov::intel_gpu::optimize_data(true));
+    network network(engine, topology, config);
+
+    auto prog = network.get_program();
+    ASSERT_NE(prog, nullptr);
+    ASSERT_TRUE(prog->get_node("shape_of").is_in_shape_of_subgraph());
+    ASSERT_FALSE(prog->get_node("random_uniform").is_in_shape_of_subgraph());
+
+    auto input_mem = engine.allocate_memory({ov::PartialShape{4, 8}, data_types::f32, format::bfyx});
+    set_values(input_mem, std::vector<float>(32, 0.f));
+    std::vector<std::vector<float>> results;
+    for (size_t i = 0; i < 3; ++i) {
+        network.set_input_data("input", input_mem);
+        auto outputs = network.execute();
+        auto output_mem = outputs.at("eltwise").get_memory();
+        cldnn::mem_lock<float, mem_lock_type::read> output_ptr(output_mem, get_test_stream());
+        results.emplace_back(output_ptr.begin(), output_ptr.end());
+    }
+    ASSERT_NE(results[0], results[1]);
+    ASSERT_NE(results[1], results[2]);
+}