From 9ee8cc38faf200e6e766f6f97dcb870173d39b56 Mon Sep 17 00:00:00 2001
From: Minseob Shin <minseob.shin11@gmail.com>
Date: Wed, 30 Sep 2026 20:56:35 +0000
Subject: [PATCH 1/7] add FP64 support for elementwise and reduce

---
 .../rocke/platform/cpp/core/dtypes.cpp        | 23 ++---
 .../rocke/platform/cpp/core/ir/core_types.cpp |  5 +
 .../rocke/platform/cpp/core/ir/ir_arith.cpp   | 11 +++
 .../rocke/platform/cpp/core/ir/ir_flow.cpp    | 10 ++
 .../rocke/platform/cpp/core/ir/ir_mem.cpp     | 21 +++-
 .../rocke/platform/cpp/core/ir/ir_tile.cpp    | 51 +++++++---
 .../platform/cpp/core/lower_hip/core.cpp      | 52 ++++++++--
 .../cpp/core/lower_hip/lower_hip_arith.cpp    | 22 ++++-
 .../platform/cpp/core/lower_llvm/arith.cpp    |  4 +-
 .../platform/cpp/core/lower_llvm/core.cpp     | 22 ++++-
 .../platform/cpp/core/lower_llvm/data.cpp     |  2 +
 .../platform/cpp/core/lower_llvm/mem.cpp      |  6 +-
 .../platform/cpp/helpers/distribution.cpp     | 10 +-
 .../rocke/platform/cpp/helpers/io.cpp         | 11 ++-
 .../rocke/platform/cpp/helpers/reduction.cpp  | 22 ++++-
 .../rocke/platform/cpp/helpers/sweep.cpp      | 12 ++-
 .../platform/cpp/helpers/tensor_view.cpp      | 30 +++++-
 .../rocke/helper_rocke.helpers.reduction.h    |  2 +-
 .../rocke/helper_rocke.helpers.tensor_view.h  | 10 ++
 .../cpp/include/rocke/instance_elementwise.h  |  5 +-
 .../cpp/include/rocke/instance_reduce.h       |  5 +-
 .../rocke/platform/cpp/include/rocke/ir.h     |  3 +
 .../cpp/include/rocke/lower_hip_internal.h    |  4 +
 .../cpp/include/rocke/lower_llvm_internal.h   |  1 +
 .../cpp/instances/common/elementwise.cpp      | 72 +++++++++++---
 .../platform/cpp/instances/common/reduce.cpp  | 81 +++++++++++----
 .../platform/dsl_docs/ir_lowering/ir_model.md |  7 +-
 .../optimization/optimization_runbook.md      |  5 +-
 .../primitives/wave_and_cross_lane.md         |  2 +-
 .../dsl_docs/reference/op_vocabulary.md       |  4 +-
 .../rocke/platform/python/rocke/__init__.py   |  2 +
 .../platform/python/rocke/core/__init__.py    |  2 +
 .../platform/python/rocke/core/dtypes.py      |  4 +
 .../rocke/platform/python/rocke/core/ir.py    | 85 +++++++++++-----
 .../platform/python/rocke/core/lower_hip.py   | 30 +++++-
 .../platform/python/rocke/core/lower_llvm.py  | 42 +++++++-
 .../python/rocke/helpers/distribution.py      | 39 +++++---
 .../rocke/platform/python/rocke/helpers/io.py | 10 +-
 .../python/rocke/helpers/reduction.py         | 24 +++--
 .../platform/python/rocke/helpers/sweep.py    |  8 +-
 .../python/rocke/helpers/tensor_view.py       | 20 +++-
 .../python/rocke/instances/SUPPORT_MATRIX.md  |  4 +
 .../rocke/instances/common/elementwise.py     | 39 ++++++--
 .../python/rocke/instances/common/reduce.py   | 40 +++++---
 .../tests/core/test_amdgpu_lowering_ops.py    |  4 +-
 .../rocke/platform/tests/core/test_dtypes.cpp |  2 +
 .../rocke/platform/tests/core/test_dtypes.py  |  3 +
 .../tests/instances/differential/numeric.py   | 99 ++++++++++++-------
 .../tests/instances/parity/elementwise_emit.c | 18 ++++
 .../instances/parity/elementwise_emit.py      |  7 ++
 .../tests/instances/parity/reduce_emit.c      | 42 +++++++-
 .../tests/instances/parity/reduce_emit.py     | 34 ++++++-
 .../rocke/platform/tests/test_rocke.py        | 41 ++++++++
 53 files changed, 898 insertions(+), 216 deletions(-)

diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/dtypes.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/dtypes.cpp
index 362d722b9e26..d3035a346b12 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/dtypes.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/dtypes.cpp
@@ -29,7 +29,8 @@ static const rocke_dtype_alias_t k_dtype_aliases[] = {
     {"f16", "fp16"},        {"half", "fp16"},         {"fp16", "fp16"},
     {"bf16", "bf16"},       {"bfloat16", "bf16"},     {"f32", "fp32"},
     {"float", "fp32"},      {"fp32", "fp32"},         {"tf32", "tf32"},
-    {"xf32", "tf32"},       {"iu8", "iu8"},           {"iu4", "iu4"},
+    {"xf32", "tf32"},       {"f64", "fp64"},          {"double", "fp64"},
+    {"fp64", "fp64"},       {"iu8", "iu8"},           {"iu4", "iu4"},
     {"i8", "i8"},           {"int8", "i8"},           {"i4", "i4"},
     {"int4", "i4"},         {"i32", ROCKE_DTYPE_I32}, {"int32", ROCKE_DTYPE_I32},
     {"fp8", "fp8e4m3"},     {"fp8e4m3", "fp8e4m3"},   {"e4m3", "fp8e4m3"},
@@ -123,16 +124,16 @@ const char* rocke_normalize_dtype(const char* name, char* lowered, size_t cap)
 }
 
 static const rocke_dtype_info_t k_dtype_info[] = {
-    {"fp32", ROCKE_DTYPE_FLOAT, 32},   {"tf32", ROCKE_DTYPE_FLOAT, 32},
-    {"fp16", ROCKE_DTYPE_FLOAT, 16},   {"bf16", ROCKE_DTYPE_FLOAT, 16},
-    {"fp8e4m3", ROCKE_DTYPE_FLOAT, 8}, {"bf8e5m2", ROCKE_DTYPE_FLOAT, 8},
-    {"fp6e2m3", ROCKE_DTYPE_FLOAT, 6}, {"fp6e3m2", ROCKE_DTYPE_FLOAT, 6},
-    {"fp4e2m1", ROCKE_DTYPE_FLOAT, 4}, {"e8m0", ROCKE_DTYPE_FLOAT, 8},
-    {"e5m3", ROCKE_DTYPE_FLOAT, 8},    {"i4", ROCKE_DTYPE_INTEGER, 4},
-    {"iu4", ROCKE_DTYPE_INTEGER, 4},   {"i8", ROCKE_DTYPE_INTEGER, 8},
-    {"iu8", ROCKE_DTYPE_INTEGER, 8},   {"i16", ROCKE_DTYPE_INTEGER, 16},
-    {"i32", ROCKE_DTYPE_INTEGER, 32},  {"i64", ROCKE_DTYPE_INTEGER, 64},
-    {"i1", ROCKE_DTYPE_PREDICATE, 1},
+    {"fp64", ROCKE_DTYPE_FLOAT, 64},   {"fp32", ROCKE_DTYPE_FLOAT, 32},
+    {"tf32", ROCKE_DTYPE_FLOAT, 32},   {"fp16", ROCKE_DTYPE_FLOAT, 16},
+    {"bf16", ROCKE_DTYPE_FLOAT, 16},   {"fp8e4m3", ROCKE_DTYPE_FLOAT, 8},
+    {"bf8e5m2", ROCKE_DTYPE_FLOAT, 8}, {"fp6e2m3", ROCKE_DTYPE_FLOAT, 6},
+    {"fp6e3m2", ROCKE_DTYPE_FLOAT, 6}, {"fp4e2m1", ROCKE_DTYPE_FLOAT, 4},
+    {"e8m0", ROCKE_DTYPE_FLOAT, 8},    {"e5m3", ROCKE_DTYPE_FLOAT, 8},
+    {"i4", ROCKE_DTYPE_INTEGER, 4},    {"iu4", ROCKE_DTYPE_INTEGER, 4},
+    {"i8", ROCKE_DTYPE_INTEGER, 8},    {"iu8", ROCKE_DTYPE_INTEGER, 8},
+    {"i16", ROCKE_DTYPE_INTEGER, 16},  {"i32", ROCKE_DTYPE_INTEGER, 32},
+    {"i64", ROCKE_DTYPE_INTEGER, 64},  {"i1", ROCKE_DTYPE_PREDICATE, 1},
 };
 
 const rocke_dtype_info_t* rocke_dtype_info(const char* name)
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/core_types.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/core_types.cpp
index 7c2230c5c50e..34eaaef3de91 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/core_types.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/core_types.cpp
@@ -54,6 +54,7 @@ ROCKE_SCALAR_SINGLETON(rocke_bf16, ROCKE_SCALAR_BF16, "bf16")
 ROCKE_SCALAR_SINGLETON(rocke_f16, ROCKE_SCALAR_F16, "f16")
 ROCKE_SCALAR_SINGLETON(rocke_f32, ROCKE_SCALAR_F32, "f32")
 ROCKE_SCALAR_SINGLETON(rocke_tf32, ROCKE_SCALAR_TF32, "tf32")
+ROCKE_SCALAR_SINGLETON(rocke_f64, ROCKE_SCALAR_F64, "f64")
 ROCKE_SCALAR_SINGLETON(rocke_fp8e4m3, ROCKE_SCALAR_FP8E4M3, "fp8e4m3")
 ROCKE_SCALAR_SINGLETON(rocke_bf8e5m2, ROCKE_SCALAR_BF8E5M2, "bf8e5m2")
 
@@ -87,6 +88,8 @@ const rocke_type_t* rocke_scalar_by_name(const char* name)
         return rocke_tf32();
     if(strcmp(name, "f32") == 0)
         return rocke_f32();
+    if(strcmp(name, "f64") == 0)
+        return rocke_f64();
     if(strcmp(name, "fp8e4m3") == 0 || strcmp(name, "e4m3") == 0)
         return rocke_fp8e4m3();
     if(strcmp(name, "bf8e5m2") == 0)
@@ -113,6 +116,8 @@ const rocke_type_t* rocke_dtype_to_ir_type(const char* dtype)
         return rocke_f16();
     if(strcmp(info->name, "fp32") == 0)
         return rocke_f32();
+    if(strcmp(info->name, "fp64") == 0)
+        return rocke_f64();
     return rocke_scalar_by_name(info->name);
 }
 
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_arith.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_arith.cpp
index 7d0b95e6cf4b..1e2f7c367ef8 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_arith.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_arith.cpp
@@ -51,6 +51,17 @@ rocke_value_t* rocke_b_const_f32(rocke_ir_builder_t* b, double value)
     return rocke_i_op1(b, ROCKE_OP_ARITH_CONSTANT, NULL, 0, rocke_f32(), &a, "c");
 }
 
+rocke_value_t* rocke_b_const_f64(rocke_ir_builder_t* b, double value)
+{
+    rocke_attr_map_t a;
+    if(!rocke_i_live(b))
+        return NULL;
+    a = rocke_i_attrs(b);
+    rocke_attr_set_float(b, &a, "value", value);
+    rocke_attr_set_str(b, &a, "ity", "f64");
+    return rocke_i_op1(b, ROCKE_OP_ARITH_CONSTANT, NULL, 0, rocke_f64(), &a, "c");
+}
+
 rocke_value_t* rocke_b_fp16_zero(rocke_ir_builder_t* b)
 {
     rocke_attr_map_t a;
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_flow.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_flow.cpp
index cbe33f0bf04c..7d9b72aa1478 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_flow.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_flow.cpp
@@ -533,6 +533,16 @@ rocke_value_t* rocke_b_warp_shuffle_xor(rocke_ir_builder_t* b, rocke_value_t* v,
         return (rocke_value_t*)rocke_i_set_err(
             b, ROCKE_ERR_VALUE, "warp_shuffle_xor: NULL operand");
 
+    /* f64: split into two i32 halves and shuffle each on the 32-bit path. */
+    if(rocke_i_type_is(v->type, "f64"))
+    {
+        rocke_value_t* halves = rocke_b_bitcast(b, v, rocke_vector_type(b, rocke_i32(), 2));
+        rocke_value_t* parts[2];
+        parts[0] = rocke_b_warp_shuffle_xor(b, rocke_b_vec_extract(b, halves, 0), lane_xor);
+        parts[1] = rocke_b_warp_shuffle_xor(b, rocke_b_vec_extract(b, halves, 1), lane_xor);
+        return rocke_b_bitcast(b, rocke_b_vec_pack(b, parts, 2, rocke_i32()), rocke_f64());
+    }
+
     if(lane_xor >= 1 && lane_xor <= 31)
     {
         /* Intra-32-lane XOR -> ds_swizzle (1 LDS op, no addr-compute). */
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_mem.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_mem.cpp
index fa16a53af676..268c913a3234 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_mem.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_mem.cpp
@@ -298,12 +298,19 @@ rocke_value_t* rocke_b_global_load_vN_ex(rocke_ir_builder_t* b,
             return (rocke_value_t*)rocke_i_set_err(
                 b, ROCKE_ERR_VALUE, "unsupported vector width for %s global_load_vN: %d", en, n);
     }
+    else if(rocke_i_type_is(dtype, "f64"))
+    {
+        elem_bytes = 8;
+        if(n != 2)
+            return (rocke_value_t*)rocke_i_set_err(
+                b, ROCKE_ERR_VALUE, "unsupported vector width for %s global_load_vN: %d", en, n);
+    }
     else
     {
         return (rocke_value_t*)rocke_i_set_err(
             b,
             ROCKE_ERR_VALUE,
-            "global_load_vN supports f16/bf16/i16/f32/i32/tf32/fp8e4m3/bf8e5m2/i8, got %s",
+            "global_load_vN supports f16/bf16/i16/f32/i32/tf32/fp8e4m3/bf8e5m2/i8/f64, got %s",
             en);
     }
     streaming = mem_opts_streaming(b, opts, "global_load_vN");
@@ -424,12 +431,22 @@ void rocke_b_global_store_vN_ex(rocke_ir_builder_t* b,
     {
         elem_bytes = 1;
     }
+    else if(rocke_i_type_is(et, "f64"))
+    {
+        elem_bytes = 8;
+        if(n > 2)
+        {
+            (void)rocke_i_set_err(
+                b, ROCKE_ERR_VALUE, "global_store_vN n=%d not supported for %s", n, en);
+            return;
+        }
+    }
     else
     {
         (void)rocke_i_set_err(
             b,
             ROCKE_ERR_VALUE,
-            "global_store_vN supports f16/bf16/i16/f32/i32/tf32/i8/fp8e4m3/bf8e5m2, got %s",
+            "global_store_vN supports f16/bf16/i16/f32/i32/tf32/i8/fp8e4m3/bf8e5m2/f64, got %s",
             en);
         return;
     }
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_tile.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_tile.cpp
index b42c30f800a3..4e7720febfda 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_tile.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_tile.cpp
@@ -87,7 +87,7 @@ static const char* rocke_mma_result_hint(const char* op_id)
     return "acc";
 }
 
-/* element-byte width helper (mirrors smem_store_vN's inline ternary). */
+/* element-byte width helper (mirrors smem_store_vN's elem_bytes). */
 static int rocke_elem_bytes_name(const char* elem_name)
 {
     if(!elem_name)
@@ -99,6 +99,10 @@ static int rocke_elem_bytes_name(const char* elem_name)
     {
         return 1;
     }
+    if(strcmp(elem_name, "f64") == 0)
+    {
+        return 8;
+    }
     if(strcmp(elem_name, "f32") == 0 || strcmp(elem_name, "i32") == 0
        || strcmp(elem_name, "tf32") == 0)
     {
@@ -240,6 +244,7 @@ void rocke_b_smem_store_vN(rocke_ir_builder_t* b,
     const char* elem_name;
     static const int allowed_8bit[] = {2, 4, 8, 16};
     static const int allowed_other[] = {2, 4, 8};
+    static const int allowed_64bit[] = {2};
     int elem_bytes;
     if(!rocke_i_live(b))
     {
@@ -271,10 +276,19 @@ void rocke_b_smem_store_vN(rocke_ir_builder_t* b,
     }
     elem_name = value->type->elem->name;
     {
-        bool eight = (strcmp(elem_name, "i8") == 0 || strcmp(elem_name, "fp8e4m3") == 0
-                      || strcmp(elem_name, "bf8e5m2") == 0);
-        const int* allowed = eight ? allowed_8bit : allowed_other;
-        int acount = eight ? 4 : 3;
+        const int* allowed = allowed_other;
+        int acount = 3;
+        if(strcmp(elem_name, "i8") == 0 || strcmp(elem_name, "fp8e4m3") == 0
+           || strcmp(elem_name, "bf8e5m2") == 0)
+        {
+            allowed = allowed_8bit;
+            acount = 4;
+        }
+        else if(strcmp(elem_name, "f64") == 0)
+        {
+            allowed = allowed_64bit;
+            acount = 1;
+        }
         if(!rocke_n_in(n, allowed, acount))
         {
             rocke_i_set_err(b,
@@ -354,6 +368,7 @@ rocke_value_t* rocke_b_smem_load_vN(rocke_ir_builder_t* b,
     static const int allowed_8bit[] = {1, 2, 4, 8, 12, 16};
     static const int allowed_16bit[] = {1, 2, 4, 6, 8};
     static const int allowed_32bit[] = {1, 2, 3, 4, 8};
+    static const int allowed_64bit[] = {1, 2};
     char hint[16];
     if(!rocke_i_live(b))
     {
@@ -366,21 +381,33 @@ rocke_value_t* rocke_b_smem_load_vN(rocke_ir_builder_t* b,
     dn = dtype->name;
     if(!(strcmp(dn, "f16") == 0 || strcmp(dn, "bf16") == 0 || strcmp(dn, "f32") == 0
          || strcmp(dn, "tf32") == 0 || strcmp(dn, "i32") == 0 || strcmp(dn, "fp8e4m3") == 0
-         || strcmp(dn, "bf8e5m2") == 0 || strcmp(dn, "i8") == 0))
+         || strcmp(dn, "bf8e5m2") == 0 || strcmp(dn, "i8") == 0 || strcmp(dn, "f64") == 0))
     {
         return (rocke_value_t*)rocke_i_set_err(
             b,
             ROCKE_ERR_VALUE,
             "smem_load_vN supports f16 / bf16 / f32 / i32 / tf32 / fp8e4m3 / "
-            "bf8e5m2 / i8, got %s",
+            "bf8e5m2 / i8 / f64, got %s",
             dn);
     }
     {
-        bool eight
-            = (strcmp(dn, "fp8e4m3") == 0 || strcmp(dn, "bf8e5m2") == 0 || strcmp(dn, "i8") == 0);
-        bool half = strcmp(dn, "f16") == 0 || strcmp(dn, "bf16") == 0;
-        const int* allowed = eight ? allowed_8bit : half ? allowed_16bit : allowed_32bit;
-        int acount = eight ? 6 : 5;
+        const int* allowed = allowed_32bit;
+        int acount = 5;
+        if(strcmp(dn, "fp8e4m3") == 0 || strcmp(dn, "bf8e5m2") == 0 || strcmp(dn, "i8") == 0)
+        {
+            allowed = allowed_8bit;
+            acount = 6;
+        }
+        else if(strcmp(dn, "f16") == 0 || strcmp(dn, "bf16") == 0)
+        {
+            allowed = allowed_16bit;
+            acount = 5;
+        }
+        else if(strcmp(dn, "f64") == 0)
+        {
+            allowed = allowed_64bit;
+            acount = 2;
+        }
         if(!rocke_n_in(n, allowed, acount))
         {
             return (rocke_value_t*)rocke_i_set_err(
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/core.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/core.cpp
index 2df181aaa4e6..1ed97aabb759 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/core.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/core.cpp
@@ -12,7 +12,7 @@
  *   - exception-based errors and the NULL guard (rocke_h_fail / rocke_h_live),
  *   - naming / type mapping (rocke_h_name / rocke_h_type_to_hip / rocke_h_hip_scalar /
  *     rocke_h_vec_prefix),
- *   - float literal formatting (rocke_h_f32_literal),
+ *   - float literal formatting (rocke_h_f32_literal / rocke_h_f64_literal),
  *   - waitcnt encoders (rocke_h_encode_waitcnt + the two raw encoders),
  *   - arch gates (rocke_h_require_ds_read_tr / rocke_h_require_wmma_arch),
  *   - the smem storage side table (rocke_h_smem_storage / rocke_h_smem_set_storage),
@@ -76,6 +76,7 @@ const char* const ROCKE_HIP_PROLOGUE
       "_ROCKE_VEC(bf16, bf16x, 8); _ROCKE_VEC(bf16, bf16x, 16);\n"
       "_ROCKE_VEC(float, f32x, 1); _ROCKE_VEC(float, f32x, 2); _ROCKE_VEC(float, f32x, 4);\n"
       "_ROCKE_VEC(float, f32x, 8); _ROCKE_VEC(float, f32x, 16);\n"
+      "_ROCKE_VEC(double, f64x, 1); _ROCKE_VEC(double, f64x, 2);\n"
       "_ROCKE_VEC(int, i32x, 1); _ROCKE_VEC(int, i32x, 2); _ROCKE_VEC(int, i32x, 3);\n"
       "_ROCKE_VEC(int, i32x, 4); _ROCKE_VEC(int, i32x, 8); _ROCKE_VEC(int, i32x, 16);\n"
       "_ROCKE_VEC(int16_t, i16x, 1); _ROCKE_VEC(int16_t, i16x, 2);\n"
@@ -316,6 +317,10 @@ const char* rocke_h_hip_scalar(const char* ir_scalar_name)
     {
         return "float";
     }
+    if(strcmp(ir_scalar_name, "f64") == 0)
+    {
+        return "double";
+    }
     if(strcmp(ir_scalar_name, "fp8e4m3") == 0)
     {
         return "fp8e4m3";
@@ -346,6 +351,10 @@ const char* rocke_h_vec_prefix(const char* ir_scalar_name, bool full_map)
             {
                 return "f32x";
             }
+            if(strcmp(ir_scalar_name, "f64") == 0)
+            {
+                return "f64x";
+            }
             if(strcmp(ir_scalar_name, "i32") == 0 || strcmp(ir_scalar_name, "tf32") == 0)
             {
                 return "i32x";
@@ -384,9 +393,9 @@ static bool rocke_h_scalar_in_vec_map(const char* name, bool full_map)
     {
         return false;
     }
-    return strcmp(name, "tf32") == 0 || strcmp(name, "f32") == 0 || strcmp(name, "i32") == 0
-           || strcmp(name, "i16") == 0 || strcmp(name, "i8") == 0 || strcmp(name, "fp8e4m3") == 0
-           || strcmp(name, "bf8e5m2") == 0;
+    return strcmp(name, "tf32") == 0 || strcmp(name, "f32") == 0 || strcmp(name, "f64") == 0
+           || strcmp(name, "i32") == 0 || strcmp(name, "i16") == 0 || strcmp(name, "i8") == 0
+           || strcmp(name, "fp8e4m3") == 0 || strcmp(name, "bf8e5m2") == 0;
 }
 
 const char* rocke_h_vec_prefix_checked(rocke_h_lowerer_t* lw,
@@ -451,9 +460,9 @@ const char* rocke_h_type_to_hip(rocke_h_lowerer_t* lw, const rocke_type_t* t)
              * through to the KeyError. Detect the listed set explicitly so an
              * unknown vector elem is an error rather than silently "f16x". */
             if(strcmp(elem, "f16") != 0 && strcmp(elem, "bf16") != 0 && strcmp(elem, "f32") != 0
-               && strcmp(elem, "tf32") != 0 && strcmp(elem, "i32") != 0 && strcmp(elem, "i16") != 0
-               && strcmp(elem, "i8") != 0 && strcmp(elem, "fp8e4m3") != 0
-               && strcmp(elem, "bf8e5m2") != 0)
+               && strcmp(elem, "tf32") != 0 && strcmp(elem, "f64") != 0 && strcmp(elem, "i32") != 0
+               && strcmp(elem, "i16") != 0 && strcmp(elem, "i8") != 0
+               && strcmp(elem, "fp8e4m3") != 0 && strcmp(elem, "bf8e5m2") != 0)
             {
                 rocke_h_fail(lw, ROCKE_ERR_KEY, "type_to_hip: unmappable vector elem '%s'", elem);
                 return "";
@@ -516,6 +525,35 @@ const char* rocke_h_f32_literal(rocke_h_lowerer_t* lw, double val)
     return out ? out : "";
 }
 
+const char* rocke_h_f64_literal(rocke_h_lowerer_t* lw, double val)
+{
+    char* out;
+    if(!lw || !lw->b)
+    {
+        return "";
+    }
+    if(isnan(val))
+    {
+        out = rocke_arena_strdup(&lw->b->arena, "((double)NAN)");
+        return out ? out : "";
+    }
+    if(isinf(val))
+    {
+        out = rocke_arena_strdup(&lw->b->arena,
+                                 val < 0 ? "((double)-INFINITY)" : "((double)INFINITY)");
+        return out ? out : "";
+    }
+    if(val == 0.0 && signbit(val))
+    {
+        out = rocke_arena_strdup(&lw->b->arena, "-0.0");
+        return out ? out : "";
+    }
+    /* %.17g round-trips every finite double. The text may differ from CPython
+     * repr (same port hazard as rocke_h_f32_literal). */
+    out = rocke_arena_printf(&lw->b->arena, "%.17g", val);
+    return out ? out : "";
+}
+
 /* ============================== waitcnt encoders ===================== */
 
 static int h_clamp(int v, int lo, int hi)
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_arith.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_arith.cpp
index eb0a283f2e4d..bf96281ef86a 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_arith.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_arith.cpp
@@ -67,6 +67,8 @@ static void h_binary(rocke_h_lowerer_t* lw, const rocke_op_t* op, const char* c_
  *             self._emit(f"{cpp_t} {_name(res)} = (fp16){literal};")
  *         else:
  *             self._emit(f"{cpp_t} {_name(res)} = {literal};")
+ *     elif ity == "f64":
+ *         self._emit(f"{cpp_t} {_name(res)} = {_f64_literal(float(val))};")
  *     else:
  *         self._emit(f"{cpp_t} {_name(res)} = {val};") */
 static rocke_status_t rocke_h_op_arith_constant(rocke_h_lowerer_t* lw, const rocke_op_t* op)
@@ -107,6 +109,13 @@ static rocke_status_t rocke_h_op_arith_constant(rocke_h_lowerer_t* lw, const roc
             rocke_h_emitf(lw, "%s %s = %s;", cpp_t, rocke_h_name(lw, res), literal);
         }
     }
+    else if(strcmp(ity, "f64") == 0)
+    {
+        double v = (val->kind == ROCKE_ATTR_FLOAT) ? val->u.f
+                   : (val->kind == ROCKE_ATTR_INT) ? (double)val->u.i
+                                                   : 0.0;
+        rocke_h_emitf(lw, "%s %s = %s;", cpp_t, rocke_h_name(lw, res), rocke_h_f64_literal(lw, v));
+    }
     else
     {
         /* integer constant: Python emits the raw `val` (an int). */
@@ -790,9 +799,20 @@ static rocke_status_t rocke_h_op_math_log2(rocke_h_lowerer_t* lw, const rocke_op
     return lw->status;
 }
 
-/* def _op_math_rcp(self, op): -- __builtin_amdgcn_rcpf, promote/demote else. */
+/* def _op_math_rcp(self, op): -- f64 divides in double, else
+ * __builtin_amdgcn_rcpf with promote/demote. */
 static rocke_status_t rocke_h_op_math_rcp(rocke_h_lowerer_t* lw, const rocke_op_t* op)
 {
+    const rocke_value_t* r = h_res(op);
+    if(r->type->name && strcmp(r->type->name, "f64") == 0)
+    {
+        rocke_h_emitf(lw,
+                      "%s %s = 1.0 / %s;",
+                      rocke_h_type_to_hip(lw, r->type),
+                      rocke_h_name(lw, r),
+                      rocke_h_name(lw, op->operands[0]));
+        return lw->status;
+    }
     h_amdgcn_unary(lw, op, "__builtin_amdgcn_rcpf");
     return lw->status;
 }
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/arith.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/arith.cpp
index 62c5c1750ad1..df6b60af0aa9 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/arith.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/arith.cpp
@@ -429,7 +429,9 @@ static void ll_fminmax(rocke_lower_t* L, const rocke_op_t* op, const char* intri
     const rocke_value_t* a = op->operands[0];
     const rocke_value_t* b = op->operands[1];
     const char* ty_name = a->type ? a->type->name : NULL;
-    const char* llvm_ty = ll_fp_llvm_ty(ty_name);
+    /* Python's fmax / fmin dicts (unlike the shared FP map) also accept f64. */
+    const char* llvm_ty
+        = (ty_name && strcmp(ty_name, "f64") == 0) ? "double" : ll_fp_llvm_ty(ty_name);
     if(llvm_ty == NULL)
     {
         rocke_ll_fail(L,
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp
index 47738b9bbcb6..9e025ce0b5c4 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp
@@ -620,6 +620,8 @@ const char* rocke_ll_llvm_type(rocke_lower_t* L, const rocke_type_t* t)
             return "i8";
         if(strcmp(n, "f32") == 0)
             return "float";
+        if(strcmp(n, "f64") == 0)
+            return "double";
     }
     rocke_ll_fail(L, ROCKE_ERR_NOTIMPL, "no LLVM mapping for type %s", n ? n : "(null)");
 }
@@ -718,7 +720,7 @@ const char* rocke_ll_smem_storage_type(rocke_lower_t* L, const rocke_type_t* sme
 }
 
 /* ====================================================================== */
-/* FP hex constants (Python _fp32_hex / _fp16_hex)                        */
+/* FP hex constants (Python _fp32_hex / _fp16_hex / _fp64_hex)            */
 /* ====================================================================== */
 
 const char* rocke_ll_fp32_hex(rocke_lower_t* L, double x)
@@ -732,6 +734,14 @@ const char* rocke_ll_fp32_hex(rocke_lower_t* L, double x)
     return rocke_arena_printf(&L->arena, "0x%016llX", (unsigned long long)bits);
 }
 
+const char* rocke_ll_fp64_hex(rocke_lower_t* L, double x)
+{
+    /* LLVM textual IR spells `double` hex constants as the raw IEEE-754 bits. */
+    uint64_t bits;
+    memcpy(&bits, &x, sizeof bits);
+    return rocke_arena_printf(&L->arena, "0x%016llX", (unsigned long long)bits);
+}
+
 const char* rocke_ll_fp16_hex(rocke_lower_t* L, double x)
 {
     /* LLVM IR: half 0xH<4 hex>. Convert double -> IEEE-754 binary16 (round to
@@ -894,6 +904,12 @@ const char* rocke_ll_operand(rocke_lower_t* L, const rocke_value_t* v)
             rocke_attr_get_float(&op->attrs, "value", &fv);
             return rocke_ll_fp16_hex(L, fv);
         }
+        if(strcmp(ity, "f64") == 0)
+        {
+            double fv = 0.0;
+            rocke_attr_get_float(&op->attrs, "value", &fv);
+            return rocke_ll_fp64_hex(L, fv);
+        }
         int64_t iv = 0;
         if(!rocke_attr_get_int(&op->attrs, "value", &iv))
         {
@@ -1203,7 +1219,7 @@ static int ll_smem_seg_size(const rocke_type_t* stype)
         eb = 2;
     else if(strcmp(n, "i32") == 0 || strcmp(n, "tf32") == 0 || strcmp(n, "f32") == 0)
         eb = 4;
-    else if(strcmp(n, "i64") == 0)
+    else if(strcmp(n, "i64") == 0 || strcmp(n, "f64") == 0)
         eb = 8;
     else
         eb = 2; /* default */
@@ -1219,6 +1235,8 @@ static int ll_smem_align(const rocke_type_t* stype)
     if(!stype || !stype->elem || !stype->elem->name)
         return 4;
     const char* n = stype->elem->name;
+    if(strcmp(n, "f64") == 0)
+        return 8;
     if(strcmp(n, "i8") == 0 || strcmp(n, "fp8e4m3") == 0 || strcmp(n, "bf8e5m2") == 0)
         return 16;
     return 4;
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/data.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/data.cpp
index be800023bb0c..8f87200c6956 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/data.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/data.cpp
@@ -153,9 +153,11 @@ const rocke_ll_decl_t ROCKE_LL_INTRINSIC_DECLS[] = {
     {"maxnum.f32", "declare float @llvm.maxnum.f32(float, float)"},
     {"maxnum.f16", "declare half @llvm.maxnum.f16(half, half)"},
     {"maxnum.bf16", "declare bfloat @llvm.maxnum.bf16(bfloat, bfloat)"},
+    {"maxnum.f64", "declare double @llvm.maxnum.f64(double, double)"},
     {"minnum.f32", "declare float @llvm.minnum.f32(float, float)"},
     {"minnum.f16", "declare half @llvm.minnum.f16(half, half)"},
     {"minnum.bf16", "declare bfloat @llvm.minnum.bf16(bfloat, bfloat)"},
+    {"minnum.f64", "declare double @llvm.minnum.f64(double, double)"},
     {"fabs.f32", "declare float @llvm.fabs.f32(float)"},
     {"fabs.f16", "declare half @llvm.fabs.f16(half)"},
     {"fabs.bf16", "declare bfloat @llvm.fabs.bf16(bfloat)"},
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp
index 114e8ff853c6..6a22dfd2ef87 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp
@@ -49,7 +49,7 @@ static int ll_elem_bytes(const char* name)
     {
         return 4;
     }
-    if(strcmp(name, "i64") == 0)
+    if(strcmp(name, "i64") == 0 || strcmp(name, "f64") == 0)
     {
         return 8;
     }
@@ -697,7 +697,7 @@ static void op_tile_smem_load_vN(rocke_lower_t* L, const rocke_op_t* op)
      * 2 (not 1). Replicate that exact dict here rather than the shared
      * ll_elem_bytes (which maps fp8/bf8 -> 1), or the fp8 down-GEMM LDS reads
      * emit `align 16` instead of the Python `align 32`.
-     *   {"i8":1,"f16":2,"bf16":2,"i32":4,"f32":4,"i64":8}.get(name, 2) */
+     *   {"i8":1,"f16":2,"bf16":2,"i32":4,"f32":4,"i64":8,"f64":8}.get(name, 2) */
     const char* en = op->results[0]->type->elem->name;
     int elem_bytes = 2;
     if(en)
@@ -714,7 +714,7 @@ static void op_tile_smem_load_vN(rocke_lower_t* L, const rocke_op_t* op)
         {
             elem_bytes = 4;
         }
-        else if(strcmp(en, "i64") == 0)
+        else if(strcmp(en, "i64") == 0 || strcmp(en, "f64") == 0)
         {
             elem_bytes = 8;
         }
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/helpers/distribution.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/helpers/distribution.cpp
index ba5a8ee5b23e..f456eb657185 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/helpers/distribution.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/helpers/distribution.cpp
@@ -1115,7 +1115,10 @@ void rocke_block_tile_reduce_sync(rocke_ir_builder_t* b,
                     = rocke_b_add(b, warp, rocke_b_const_i32(b, i * num_reduce_warps));
                 rocke_value_t* idx[1];
                 idx[0] = slot;
-                rocke_b_smem_store_vN_f32(b, lds_buf, idx, 1, v, 1);
+                if(strcmp(v->type->name, "f64") == 0)
+                    rocke_b_smem_store_vN(b, lds_buf, idx, 1, v, 1);
+                else
+                    rocke_b_smem_store_vN_f32(b, lds_buf, idx, 1, v, 1);
             }
             rocke_b_region_leave(b);
         }
@@ -1135,7 +1138,10 @@ void rocke_block_tile_reduce_sync(rocke_ir_builder_t* b,
                 rocke_value_t* sidx[1];
                 rocke_value_t* vec;
                 sidx[0] = slot;
-                vec = rocke_b_smem_load_vN_f32(b, lds_buf, sidx, 1, 1);
+                if(strcmp(reduced->storage[i]->type->name, "f64") == 0)
+                    vec = rocke_b_smem_load_vN(b, lds_buf, sidx, 1, rocke_f64(), 1);
+                else
+                    vec = rocke_b_smem_load_vN_f32(b, lds_buf, sidx, 1, 1);
                 parts[idx_i] = rocke_b_vec_extract(b, vec, 0);
             }
             /* Pairwise power-of-two tree fold. */
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/helpers/io.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/helpers/io.cpp
index e90e86ded9e2..d6bcf65e178c 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/helpers/io.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/helpers/io.cpp
@@ -11,14 +11,16 @@
  *     def io_ir_type(dtype: str) -> Type:
  *         if dtype in ("f16", "fp16"): return F16
  *         if dtype == "bf16":          return BF16
+ *         if dtype == "f64":           return F64
  *         raise ValueError(
- *             f"unsupported I/O dtype {dtype!r}; expected f16/fp16/bf16")
+ *             f"unsupported I/O dtype {dtype!r}; expected f16/fp16/bf16/f64")
  *
  * Mapping invariants (must stay byte-identical to the Python so downstream IR is
  * identical):
  *   "f16"  -> rocke_f16()    (Python F16)
  *   "fp16" -> rocke_f16()    (alias; Python F16)
  *   "bf16" -> rocke_bf16()   (Python BF16)
+ *   "f64"  -> rocke_f64()    (Python F64)
  *   else   -> NULL         (Python ValueError)
  */
 
@@ -52,6 +54,11 @@ const rocke_type_t* rocke_io_ir_type(const char* dtype)
     {
         return rocke_bf16();
     }
+    /* `dtype == "f64"` -> F64 */
+    if(strcmp(dtype, "f64") == 0)
+    {
+        return rocke_f64();
+    }
     /* Python: raise ValueError. No builder here, so signal via NULL. */
     return NULL;
 }
@@ -75,7 +82,7 @@ const rocke_type_t* rocke_b_io_ir_type(rocke_ir_builder_t* b, const char* dtype)
         return (const rocke_type_t*)rocke_i_set_err(
             b,
             ROCKE_ERR_VALUE,
-            "unsupported I/O dtype %s%s%s; expected f16/fp16/bf16",
+            "unsupported I/O dtype %s%s%s; expected f16/fp16/bf16/f64",
             dtype ? "'" : "",
             dtype ? dtype : "None",
             dtype ? "'" : "");
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/helpers/reduction.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/helpers/reduction.cpp
index 409b26f03386..c895481c9557 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/helpers/reduction.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/helpers/reduction.cpp
@@ -55,6 +55,12 @@ static bool rocke_red_is_f32(const rocke_value_t* v)
            && strcmp(v->type->name, "f32") == 0;
 }
 
+static bool rocke_red_is_f64(const rocke_value_t* v)
+{
+    return v != NULL && v->type != NULL && v->type->name != NULL
+           && strcmp(v->type->name, "f64") == 0;
+}
+
 static bool rocke_red_is_i32(const rocke_value_t* v)
 {
     return v != NULL && v->type != NULL && v->type->name != NULL
@@ -381,13 +387,14 @@ rocke_value_t* rocke_block_lds_reduce_with_wave_prologue(rocke_ir_builder_t* b,
     rocke_value_t* warp;
     rocke_value_t** parts;
     int w;
+    bool is_f64;
 
-    if(!rocke_red_is_f32(val))
+    if(!rocke_red_is_f32(val) && !rocke_red_is_f64(val))
     {
         return (rocke_value_t*)rocke_red_set_err(
             b,
             ROCKE_ERR_VALUE,
-            "block_lds_reduce_with_wave_prologue expects f32 input, got %s",
+            "block_lds_reduce_with_wave_prologue expects f32 or f64 input, got %s",
             (val && val->type) ? val->type->name : "<null>");
     }
 
@@ -402,13 +409,17 @@ rocke_value_t* rocke_block_lds_reduce_with_wave_prologue(rocke_ir_builder_t* b,
     c_wave = rocke_b_const_i32(b, wave_size);
     lane = rocke_b_mod(b, tid, c_wave);
     warp = rocke_b_div(b, tid, c_wave);
+    is_f64 = rocke_red_is_f64(val);
     {
         rocke_if_t iff = rocke_b_scf_if(b, rocke_b_cmp_eq(b, lane, rocke_b_const_i32(b, 0)));
         rocke_b_region_enter(b, iff.then_region);
         {
             rocke_value_t* idx[1];
             idx[0] = warp;
-            rocke_b_smem_store_vN_f32(b, lds_buf, idx, 1, warp_partial, 1);
+            if(is_f64)
+                rocke_b_smem_store_vN(b, lds_buf, idx, 1, warp_partial, 1);
+            else
+                rocke_b_smem_store_vN_f32(b, lds_buf, idx, 1, warp_partial, 1);
         }
         rocke_b_region_leave(b);
     }
@@ -425,7 +436,10 @@ rocke_value_t* rocke_block_lds_reduce_with_wave_prologue(rocke_ir_builder_t* b,
         rocke_value_t* idx[1];
         rocke_value_t* v_vec;
         idx[0] = rocke_b_const_i32(b, w);
-        v_vec = rocke_b_smem_load_vN_f32(b, lds_buf, idx, 1, 1);
+        if(is_f64)
+            v_vec = rocke_b_smem_load_vN(b, lds_buf, idx, 1, rocke_f64(), 1);
+        else
+            v_vec = rocke_b_smem_load_vN_f32(b, lds_buf, idx, 1, 1);
         parts[w] = rocke_b_vec_extract(b, v_vec, 0);
     }
     return rocke_red_tree_reduce_scalars(b, combine, parts, num_warps);
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/helpers/sweep.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/helpers/sweep.cpp
index 5aba21f52ada..0acd61de0a79 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/helpers/sweep.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/helpers/sweep.cpp
@@ -18,6 +18,8 @@
 #include "rocke/arena.h"
 #include "rocke/ir_internal.h" /* rocke_i_set_err for Python-ValueError parity */
 
+#include <string.h>
+
 /* ------------------------------------------------------------------------
  * sweep_row_chunks
  *
@@ -34,7 +36,10 @@
  *   for k in range(chunks_per_thread):
  *       n_off = b.add(b.mul(b.const_i32(k * block_size), c_vec),
  *                     b.mul(tid, c_vec))
- *       x_scalars = tile.load_vec_as_f32(b, b.const_i32(0), n_off, n=vec)
+ *       if tile.dtype.name == "f64":
+ *           x_scalars = tile.load_vec_scalars(b, b.const_i32(0), n_off, n=vec)
+ *       else:
+ *           x_scalars = tile.load_vec_as_f32(b, b.const_i32(0), n_off, n=vec)
  *       if cache: cached.extend(x_scalars)
  *       if body is not None: body(n_off, x_scalars)
  *   return RowChunkSweepResult(cached=cached, chunks_per_thread=...)
@@ -130,7 +135,10 @@ rocke_row_chunk_sweep_result_t rocke_sweep_row_chunks(rocke_ir_builder_t* b,
          * which spells the literal-0 column index inline at the load. */
         local_indices[0] = rocke_b_const_i32(b, 0);
         local_indices[1] = n_off;
-        rocke_tile_window_load_vec_as_f32(b, t, local_indices, 2, vec, scratch);
+        if(strcmp(rocke_tile_window_dtype(t)->name, "f64") == 0)
+            rocke_tile_window_load_vec_scalars(b, t, local_indices, 2, vec, scratch);
+        else
+            rocke_tile_window_load_vec_as_f32(b, t, local_indices, 2, vec, scratch);
 
         if(cache && res.cached != NULL)
         {
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/helpers/tensor_view.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/helpers/tensor_view.cpp
index fc9362193d1e..59371c288efc 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/helpers/tensor_view.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/helpers/tensor_view.cpp
@@ -311,7 +311,7 @@ rocke_value_t* rocke_tensor_view_load_vec(rocke_ir_builder_t* b,
         return NULL;
     }
     off = rocke_tensor_descriptor_offset(b, &v->desc, indices, num_indices);
-    if(rocke_tv_name_is(dt, "f16") || rocke_tv_name_is(dt, "bf16"))
+    if(rocke_tv_name_is(dt, "f16") || rocke_tv_name_is(dt, "bf16") || rocke_tv_name_is(dt, "f64"))
         return rocke_b_global_load_vN(b, v->base, off, dt, n, 0);
     if(rocke_tv_name_is(dt, "f32"))
     {
@@ -511,6 +511,34 @@ void rocke_tile_window_store_vec(rocke_ir_builder_t* b,
     rocke_tensor_view_store_vec(b, w->view, gidx, w->rank, value, n);
 }
 
+void rocke_tile_window_load_vec_scalars(rocke_ir_builder_t* b,
+                                        const rocke_tile_window_t* w,
+                                        rocke_value_t* const* local_indices,
+                                        int num_indices,
+                                        int n,
+                                        rocke_value_t** out)
+{
+    /* Python TileWindow.load_vec_scalars:
+     *   if n == 1:
+     *       return [self.load_scalar(b, *local_indices)]
+     *   v = self.load_vec(b, *local_indices, n=n)
+     *   return [b.vec_extract(v, i) for i in range(n)]
+     */
+    int i;
+    if(b == NULL || w == NULL || w->view == NULL || out == NULL)
+        return;
+    if(n == 1)
+    {
+        out[0] = rocke_tile_window_load_scalar(b, w, local_indices, num_indices);
+        return;
+    }
+    {
+        rocke_value_t* v = rocke_tile_window_load_vec(b, w, local_indices, num_indices, n);
+        for(i = 0; i < n; ++i)
+            out[i] = rocke_b_vec_extract(b, v, i);
+    }
+}
+
 rocke_value_t* rocke_tile_window_load_scalar(rocke_ir_builder_t* b,
                                              const rocke_tile_window_t* w,
                                              rocke_value_t* const* local_indices,
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/helper_rocke.helpers.reduction.h b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/helper_rocke.helpers.reduction.h
index c370c1ac7116..d3161929f824 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/helper_rocke.helpers.reduction.h
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/helper_rocke.helpers.reduction.h
@@ -126,7 +126,7 @@ int rocke_block_lds_reduce_pair(rocke_ir_builder_t* b,
 /* block_lds_reduce_with_wave_prologue(...) analogue: wave-XOR butterfly +
  * cross-warp LDS. Six (for wave_size=64) cross-lane shuffle stages with no LDS,
  * then one sync over a num_warps-slot scratch in ``lds_buf``. Returns NULL (and
- * sets sticky error) on non-f32 input. */
+ * sets sticky error) on input that is not f32 or f64. */
 rocke_value_t* rocke_block_lds_reduce_with_wave_prologue(rocke_ir_builder_t* b,
                                                          rocke_value_t* val,
                                                          rocke_value_t* lds_buf,
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/helper_rocke.helpers.tensor_view.h b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/helper_rocke.helpers.tensor_view.h
index 7536f1955501..3b56bb0f6a6a 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/helper_rocke.helpers.tensor_view.h
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/helper_rocke.helpers.tensor_view.h
@@ -225,6 +225,16 @@ void rocke_tile_window_store_vec(rocke_ir_builder_t* b,
                                  rocke_value_t* value,
                                  int n);
 
+/* TileWindow.load_vec_scalars(b, *local_indices, n): vector load split into
+ * ``n`` scalars in the window's own dtype (no f32 promotion). Writes them to
+ * out[0..n) (caller-provided, length >= n). */
+void rocke_tile_window_load_vec_scalars(rocke_ir_builder_t* b,
+                                        const rocke_tile_window_t* w,
+                                        rocke_value_t* const* local_indices,
+                                        int num_indices,
+                                        int n,
+                                        rocke_value_t** out);
+
 /* TileWindow.load_scalar(b, *local_indices). */
 rocke_value_t* rocke_tile_window_load_scalar(rocke_ir_builder_t* b,
                                              const rocke_tile_window_t* w,
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_elementwise.h b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_elementwise.h
index 2babf90be340..7fc8a405be28 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_elementwise.h
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_elementwise.h
@@ -7,7 +7,8 @@
  *
  * Emits a single AMDGPU kernel that walks one contiguous N-element tensor with
  * vectorised global loads/stores and applies a fused unary or binary operation
- * per element. Compute is f32 internally; I/O is f16 or bf16.
+ * per element. Compute is f32 internally for f16 or bf16 I/O; f64 I/O computes
+ * natively (exact ops only, vec=2).
  *
  *   Python (elementwise.py)               C99 (this header)
  *   -----------------------------------   --------------------------------------
@@ -53,7 +54,7 @@ extern "C" {
 /* --------------------------------------------------------- ElementwiseSpec */
 
 /* Mirror of Python ElementwiseSpec. `op` is one of the unary or binary op
- * spellings (see the dispatch tables below); `dtype` is "f16" or "bf16". */
+ * spellings (see the dispatch tables below); `dtype` is "f16", "bf16" or "f64". */
 typedef struct rocke_elementwise_spec
 {
     const char* op; /* required (no default in Python)                  */
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_reduce.h b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_reduce.h
index 591dba4348b3..d6cee4fde623 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_reduce.h
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_reduce.h
@@ -62,7 +62,7 @@ extern "C" {
  *     op: ReduceOp = "sum"          # "sum"|"max"|"min"|"mean"|"prod"
  *     block_size: int = 256
  *     vec: int = 4
- *     dtype: DType = "f16"          # "f16"|"bf16"
+ *     dtype: DType = "f16"          # "f16"|"bf16"|"f64"
  *     wave_size: int = 64
  *     name: str = "rocke_reduce2d"
  */
@@ -100,7 +100,8 @@ rocke_status_t
  *
  * Gate (mirrors reduce.is_valid_spec):
  *   - op in ("sum","max","min","mean","prod")
- *   - validate_io(IOSpecRule(dtype, block_size, vec, n_per_block))
+ *   - validate_io(IOSpecRule(dtype, block_size, vec, n_per_block)); f64 allows
+ *     only vec=2
  *
  * On reject, `reason` (if non-NULL, capacity reason_cap) receives the structured
  * message and false is returned. On accept returns true and writes "ok". */
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/ir.h b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/ir.h
index ad64e60191fb..9589acfb042e 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/ir.h
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/ir.h
@@ -103,6 +103,7 @@ typedef enum rocke_scalar_kind
     ROCKE_SCALAR_BF16,
     ROCKE_SCALAR_F16,
     ROCKE_SCALAR_F32,
+    ROCKE_SCALAR_F64,
     ROCKE_SCALAR_FP8E4M3,
     ROCKE_SCALAR_BF8E5M2,
     ROCKE_SCALAR_FP4E2M1,
@@ -620,6 +621,7 @@ const rocke_type_t* rocke_bf16(void);
 const rocke_type_t* rocke_f16(void);
 const rocke_type_t* rocke_f32(void);
 const rocke_type_t* rocke_tf32(void);
+const rocke_type_t* rocke_f64(void);
 const rocke_type_t* rocke_fp8e4m3(void);
 const rocke_type_t* rocke_bf8e5m2(void);
 const rocke_type_t* rocke_fp4e2m1(void);
@@ -737,6 +739,7 @@ rocke_value_t* rocke_b_get_param(rocke_ir_builder_t* b, const char* name);
 rocke_value_t* rocke_b_const_i32(rocke_ir_builder_t* b, int64_t value);
 rocke_value_t* rocke_b_const_i64(rocke_ir_builder_t* b, int64_t value);
 rocke_value_t* rocke_b_const_f32(rocke_ir_builder_t* b, double value);
+rocke_value_t* rocke_b_const_f64(rocke_ir_builder_t* b, double value);
 rocke_value_t* rocke_b_fp16_zero(rocke_ir_builder_t* b);
 rocke_value_t* rocke_b_zero_vec_f32(rocke_ir_builder_t* b, int n);
 rocke_value_t* rocke_b_zero_vec(rocke_ir_builder_t* b, const rocke_type_t* elem, int n);
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_hip_internal.h b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_hip_internal.h
index 0f70ba934ba6..ae2283f8c2c1 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_hip_internal.h
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_hip_internal.h
@@ -181,6 +181,10 @@ const char* rocke_h_vec_prefix_checked(rocke_h_lowerer_t* lw,
  * arena-owned string. */
 const char* rocke_h_f32_literal(rocke_h_lowerer_t* lw, double val);
 
+/* Python _f64_literal(val): C++ double-literal text for a double, special-casing
+ * nan / +/-inf like rocke_h_f32_literal. Returns an arena-owned string. */
+const char* rocke_h_f64_literal(rocke_h_lowerer_t* lw, double val);
+
 /* Python _encode_waitcnt: arch-aware s_waitcnt immediate. Selects the gfx9/10
  * split-VMCNT layout or the gfx11 contiguous layout off lw->arch. -1 on any
  * counter means "no wait" (encoded as that field's architectural max). */
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_llvm_internal.h b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_llvm_internal.h
index e130cfafe700..75280cfe9f17 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_llvm_internal.h
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_llvm_internal.h
@@ -569,6 +569,7 @@ int64_t rocke_ll_eval_constant(rocke_lower_t* L, const rocke_value_t* v);
  * Arena-owned. */
 const char* rocke_ll_fp32_hex(rocke_lower_t* L, double x);
 const char* rocke_ll_fp16_hex(rocke_lower_t* L, double x);
+const char* rocke_ll_fp64_hex(rocke_lower_t* L, double x);
 
 /* Escape a string for an LLVM asm/string literal (Python
  * _escape_llvm_asm_string): printable ASCII verbatim, else \XX hex. */
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/elementwise.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/elementwise.cpp
index d3c0376b27d8..c11fd5b1389c 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/elementwise.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/elementwise.cpp
@@ -139,6 +139,15 @@ static bool rocke_ew_reason(char* reason, size_t cap, const char* msg)
     return false;
 }
 
+/* _F64_OPS: ops computed exactly in the I/O type; the exp2-based activations
+ * are f32-only, so f64 is limited to this set. */
+static bool rocke_ew_is_f64_op(const char* op)
+{
+    return rocke_ew_streq(op, "copy") || rocke_ew_streq(op, "neg") || rocke_ew_streq(op, "abs")
+           || rocke_ew_streq(op, "relu") || rocke_ew_streq(op, "add") || rocke_ew_streq(op, "sub")
+           || rocke_ew_streq(op, "mul") || rocke_ew_streq(op, "max") || rocke_ew_streq(op, "min");
+}
+
 bool rocke_elementwise_is_valid_spec(const rocke_elementwise_spec_t* spec,
                                      char* reason,
                                      size_t reason_cap)
@@ -153,11 +162,17 @@ bool rocke_elementwise_is_valid_spec(const rocke_elementwise_spec_t* spec,
         snprintf(buf, sizeof(buf), "unknown op %s", spec->op ? spec->op : "(null)");
         return rocke_ew_reason(reason, reason_cap, buf);
     }
-    if(!(rocke_ew_streq(spec->dtype, "f16") || rocke_ew_streq(spec->dtype, "bf16")))
+    if(!(rocke_ew_streq(spec->dtype, "f16") || rocke_ew_streq(spec->dtype, "bf16")
+         || rocke_ew_streq(spec->dtype, "f64")))
     {
         snprintf(buf, sizeof(buf), "unsupported dtype %s", spec->dtype ? spec->dtype : "(null)");
         return rocke_ew_reason(reason, reason_cap, buf);
     }
+    if(rocke_ew_streq(spec->dtype, "f64") && !rocke_ew_is_f64_op(spec->op))
+    {
+        snprintf(buf, sizeof(buf), "op '%s' not supported for f64", spec->op);
+        return rocke_ew_reason(reason, reason_cap, buf);
+    }
     if(!(spec->block_size == 64 || spec->block_size == 128 || spec->block_size == 256
          || spec->block_size == 512 || spec->block_size == 1024))
     {
@@ -165,7 +180,15 @@ bool rocke_elementwise_is_valid_spec(const rocke_elementwise_spec_t* spec,
             buf, sizeof(buf), "block_size %d not in {64, 128, 256, 512, 1024}", spec->block_size);
         return rocke_ew_reason(reason, reason_cap, buf);
     }
-    if(!(spec->vec == 2 || spec->vec == 4 || spec->vec == 8))
+    if(rocke_ew_streq(spec->dtype, "f64"))
+    {
+        if(spec->vec != 2)
+        {
+            snprintf(buf, sizeof(buf), "vec %d not in {2} for f64", spec->vec);
+            return rocke_ew_reason(reason, reason_cap, buf);
+        }
+    }
+    else if(!(spec->vec == 2 || spec->vec == 4 || spec->vec == 8))
     {
         snprintf(buf, sizeof(buf), "vec %d not in {2, 4, 8}", spec->vec);
         return rocke_ew_reason(reason, reason_cap, buf);
@@ -222,7 +245,9 @@ static rocke_value_t* rocke_ew_apply_unary(rocke_ir_builder_t* b, rocke_value_t*
     }
     if(rocke_ew_streq(op, "relu"))
     {
-        return rocke_b_fmax(b, x, rocke_b_const_f32(b, 0.0));
+        rocke_value_t* zero = strcmp(x->type->name, "f64") == 0 ? rocke_b_const_f64(b, 0.0)
+                                                                : rocke_b_const_f32(b, 0.0);
+        return rocke_b_fmax(b, x, zero);
     }
     if(rocke_ew_streq(op, "exp2"))
     {
@@ -296,19 +321,34 @@ static rocke_value_t*
     return NULL;
 }
 
+/* _to_compute: promote an I/O scalar to the compute type (f64 stays native). */
+static rocke_value_t* rocke_ew_to_compute(rocke_ir_builder_t* b, rocke_value_t* x)
+{
+    return strcmp(x->type->name, "f64") == 0 ? x : rocke_b_cast_to_f32(b, x);
+}
+
+/* _from_compute: demote a compute scalar back to the I/O type (f64 stays native). */
+static rocke_value_t*
+    rocke_ew_from_compute(rocke_ir_builder_t* b, rocke_value_t* x, const rocke_type_t* io_ty)
+{
+    return strcmp(io_ty->name, "f64") == 0 ? x : rocke_b_cast_f32_to(b, x, io_ty);
+}
+
 /* ===================================================================== *
  *  Distribution-driven load / store specialised for the elementwise tile.
  *
  *  Reproduces the Python load_tile / store_tile builder-call order for the
  *  fixed single-Y, single-P, single-X distribution this instance uses. The
- *  per-thread register tile holds exactly `vec` f32 scalars (storage[0..vec)).
+ *  per-thread register tile holds exactly `vec` compute scalars (f32, or
+ *  native f64 for f64 I/O) in storage[0..vec).
  * ===================================================================== */
 
 /* load_tile(window, distribution, ps=[[tid]]) for the elementwise distribution.
  *
  *   traits.iterate_accesses() yields one base y_base = (0,)
  *   x_coords = distribution.calculate_x(ys=[const_i32(0)], ps=[[tid]])
- *   scalars  = window.load_vec_as_f32(*x_coords, n=vec)
+ *   scalars  = window.load_vec_scalars(*x_coords, n=vec)   if f64
+ *              window.load_vec_as_f32(*x_coords, n=vec)    otherwise
  *   dt[k] = scalars[k]
  *
  * `out_storage` must have capacity >= vec. Returns 1 on success, 0 on failure
@@ -340,6 +380,12 @@ static int rocke_ew_load_tile(rocke_ir_builder_t* b,
         return 0;
     }
 
+    if(strcmp(rocke_tile_window_dtype(window)->name, "f64") == 0)
+    {
+        rocke_tile_window_load_vec_scalars(b, window, x_coords, 1, vec, out_storage);
+        return rocke_ir_builder_ok(b) ? 1 : 0;
+    }
+
     /* window.load_vec_as_f32(*x_coords, n=vec):
      *   v = load_vec(x_coords, n=vec)
      *   scalars[k] = cast_to_f32(vec_extract(v, k))
@@ -356,7 +402,8 @@ static int rocke_ew_load_tile(rocke_ir_builder_t* b,
  *
  *   x_coords = calculate_x(ys=[const_i32(0)], ps=[[tid]])
  *   scalars  = storage[0..vec)
- *   window.store_vec_from_f32(*x_coords, values=scalars):
+ *   f64: packed = vec_pack(scalars, dtype); store_vec(x_coords, packed, n=vec)
+ *   otherwise window.store_vec_from_f32(*x_coords, values=scalars):
  *       casts[k] = cast_f32_to(scalars[k], dtype)
  *       packed   = vec_pack(casts, dtype)
  *       store_vec(x_coords, packed, n=vec)
@@ -391,7 +438,8 @@ static void rocke_ew_store_tile(rocke_ir_builder_t* b,
     dtype = rocke_tile_window_dtype(window);
     for(k = 0; k < vec; ++k)
     {
-        casts[k] = rocke_b_cast_f32_to(b, storage[k], dtype);
+        casts[k] = strcmp(dtype->name, "f64") == 0 ? storage[k]
+                                                   : rocke_b_cast_f32_to(b, storage[k], dtype);
     }
     packed = rocke_b_vec_pack(b, casts, vec, dtype);
     rocke_tile_window_store_vec(b, window, x_coords, 1, packed, vec);
@@ -658,12 +706,12 @@ rocke_kernel_def_t* rocke_build_elementwise(rocke_ir_builder_t* b,
                     rocke_value_t* a_s;
                     rocke_value_t* r;
                     indices[0] = idx;
-                    /* a = cast_to_f32(a_view.load_scalar([idx])) */
-                    a_s = rocke_b_cast_to_f32(
+                    /* a = _to_compute(a_view.load_scalar([idx])) */
+                    a_s = rocke_ew_to_compute(
                         b, rocke_tensor_view_load_scalar(b, &a_view, indices, 1));
                     if(is_binary)
                     {
-                        rocke_value_t* bv = rocke_b_cast_to_f32(
+                        rocke_value_t* bv = rocke_ew_to_compute(
                             b, rocke_tensor_view_load_scalar(b, &b_view, indices, 1));
                         r = rocke_ew_apply_binary(b, a_s, bv, spec->op);
                     }
@@ -671,9 +719,9 @@ rocke_kernel_def_t* rocke_build_elementwise(rocke_ir_builder_t* b,
                     {
                         r = rocke_ew_apply_unary(b, a_s, spec->op);
                     }
-                    /* c_view.store_scalar([idx], cast_f32_to(r, io_ty)) */
+                    /* c_view.store_scalar([idx], _from_compute(r, io_ty)) */
                     rocke_tensor_view_store_scalar(
-                        b, &c_view, indices, 1, rocke_b_cast_f32_to(b, r, io_ty), 0);
+                        b, &c_view, indices, 1, rocke_ew_from_compute(b, r, io_ty), 0);
                 }
                 rocke_b_region_leave(b);
             }
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/reduce.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/reduce.cpp
index a06e9d023588..1dc317c368c7 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/reduce.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/reduce.cpp
@@ -6,15 +6,18 @@
  * Row-wise reduction kernel (M,N) -> (M,1): sum/max/min/mean/prod per row.
  * Byte-identical builder-call sequence vs the Python build_reduce2d.
  *
- * Stage 1 (thread): tree_reduce f32 fold of each vec-chunk, joined onto acc.
+ * Stage 1 (thread): tree_reduce fold of each vec-chunk in the compute type
+ *   (f32, or f64 for f64 input), joined onto acc.
  * Stage 2 (cross-thread): wave-aligned sum/max -> reduce-distribution +
  *   block_tile_reduce_sync; wave-aligned min/prod -> wave-XOR prologue;
  *   non-wave -> full LDS tree.
- * Stage 3 (lane 0): scalar store of the f32 result demoted to dtype.
+ * Stage 3 (lane 0): scalar store of the result demoted to dtype (f64 is
+ *   stored as is).
  */
 
 #include "rocke/instance_reduce.h"
 
+#include <math.h>
 #include <stdio.h>
 #include <string.h>
 
@@ -33,10 +36,12 @@
  * port (helper_rocke.helpers.distribution.{h,c}), included above. */
 
 /* ===================================================================== *
- *  f32 identity constants (reduce.py module-level).
+ *  f32 / f64 identity constants (reduce.py module-level).
  * ===================================================================== */
 #define ROCKE_REDUCE_NEG_INF_F32 (-3.4028234663852886e38)
 #define ROCKE_REDUCE_POS_INF_F32 (3.4028234663852886e38)
+#define ROCKE_REDUCE_NEG_INF_F64 (-INFINITY)
+#define ROCKE_REDUCE_POS_INF_F64 (INFINITY)
 
 /* ===================================================================== *
  *  Spec defaults / properties.
@@ -111,8 +116,14 @@ rocke_status_t
 
 /* ===================================================================== *
  *  is_valid_spec(spec):
- *    op in (...); validate_io(IOSpecRule(dtype, block_size, vec, n_per_block))
+ *    op in (...); validate_io(IOSpecRule(dtype, block_size, vec, n_per_block,
+ *        allowed_dtypes=("f16","fp16","bf16","f64"),
+ *        allowed_vecs=(2,) if dtype == "f64" else (2, 4, 8)))
  * ===================================================================== */
+static const char* const rocke_reduce2d_dtypes[] = {"f16", "fp16", "bf16", "f64"};
+static const int rocke_reduce2d_vecs[] = {2, 4, 8};
+static const int rocke_reduce2d_vecs_f64[] = {2};
+
 static int rocke_reduce2d_op_ok(const char* op)
 {
     return op != NULL
@@ -148,6 +159,18 @@ bool rocke_reduce2d_is_valid_spec(const rocke_reduce2d_spec_t* spec,
     rocke_io_spec_rule_init(&rule, spec->dtype, spec->block_size, spec->vec);
     rule.n_per_block_set = 1;
     rule.n_per_block = spec->n_per_block;
+    rule.allowed_dtypes = rocke_reduce2d_dtypes;
+    rule.num_allowed_dtypes = 4;
+    if(spec->dtype != NULL && strcmp(spec->dtype, "f64") == 0)
+    {
+        rule.allowed_vecs = rocke_reduce2d_vecs_f64;
+        rule.num_allowed_vecs = 1;
+    }
+    else
+    {
+        rule.allowed_vecs = rocke_reduce2d_vecs;
+        rule.num_allowed_vecs = 3;
+    }
 
     rocke_arena_init(&arena, 0);
     ok = rocke_validate_io(&arena, &rule, &why);
@@ -159,8 +182,14 @@ bool rocke_reduce2d_is_valid_spec(const rocke_reduce2d_spec_t* spec,
     return ok != 0;
 }
 
+/* const_c = b.const_f64 if is_f64 else b.const_f32 */
+static rocke_value_t* rocke_reduce2d_const_c(rocke_ir_builder_t* b, bool is_f64, double value)
+{
+    return is_f64 ? rocke_b_const_f64(b, value) : rocke_b_const_f32(b, value);
+}
+
 /* ===================================================================== *
- *  _combine_scalar(b, combine, a, c): one f32 reduction step.
+ *  _combine_scalar(b, combine, a, c): one reduction step in the compute type.
  *  combine here is a rocke_reduce_combine_t (sum/max/min/prod).
  * ===================================================================== */
 static rocke_value_t* rocke_reduce2d_combine_scalar(rocke_ir_builder_t* b,
@@ -326,6 +355,8 @@ rocke_kernel_def_t*
         rocke_tensor_view_t x_view;
         rocke_tile_window_t x_tile;
         rocke_value_t* lds;
+        bool is_f64;
+        const rocke_type_t* compute_ty;
 
         rocke_reduce_combine_t combine;
         rocke_value_t* acc;
@@ -421,33 +452,40 @@ rocke_kernel_def_t*
             }
         }
 
-        /* lds = make_lds_view(b, dtype=F32, shape=(BS,), name_hint="lds_red").base
-         *   == smem_alloc(F32, [BS], "lds_red"); .base is the smem token. */
+        /* f16/bf16 accumulate in f32; f64 accumulates natively.
+         * is_f64 = io_ty == F64; compute_ty = F64 if is_f64 else F32 */
+        is_f64 = (io_ty == rocke_f64());
+        compute_ty = is_f64 ? rocke_f64() : rocke_f32();
+
+        /* lds = make_lds_view(b, dtype=compute_ty, shape=(BS,), name_hint="lds_red").base
+         *   == smem_alloc(compute_ty, [BS], "lds_red"); .base is the smem token. */
         {
             int lds_shape[1];
             lds_shape[0] = BS;
-            lds = rocke_b_smem_alloc(b, rocke_f32(), lds_shape, 1, "lds_red");
+            lds = rocke_b_smem_alloc(b, compute_ty, lds_shape, 1, "lds_red");
         }
 
-        /* Pick f32 identity + combiner per op. */
+        /* Pick the identity element + combiner per op (const_c = const_f64 or const_f32). */
         if(strcmp(spec->op, "sum") == 0 || strcmp(spec->op, "mean") == 0)
         {
-            acc = rocke_b_const_f32(b, 0.0);
+            acc = rocke_reduce2d_const_c(b, is_f64, 0.0);
             combine = ROCKE_REDUCE_SUM;
         }
         else if(strcmp(spec->op, "max") == 0)
         {
-            acc = rocke_b_const_f32(b, ROCKE_REDUCE_NEG_INF_F32);
+            acc = rocke_reduce2d_const_c(
+                b, is_f64, is_f64 ? ROCKE_REDUCE_NEG_INF_F64 : ROCKE_REDUCE_NEG_INF_F32);
             combine = ROCKE_REDUCE_MAX;
         }
         else if(strcmp(spec->op, "min") == 0)
         {
-            acc = rocke_b_const_f32(b, ROCKE_REDUCE_POS_INF_F32);
+            acc = rocke_reduce2d_const_c(
+                b, is_f64, is_f64 ? ROCKE_REDUCE_POS_INF_F64 : ROCKE_REDUCE_POS_INF_F32);
             combine = ROCKE_REDUCE_MIN;
         }
         else if(strcmp(spec->op, "prod") == 0)
         {
-            acc = rocke_b_const_f32(b, 1.0);
+            acc = rocke_reduce2d_const_c(b, is_f64, 1.0);
             combine = ROCKE_REDUCE_PROD;
         }
         else
@@ -478,7 +516,7 @@ rocke_kernel_def_t*
            && (combine == ROCKE_REDUCE_SUM || combine == ROCKE_REDUCE_MAX))
         {
             /* red_dist = _make_row_reduce_distribution(spec)
-             * reduced  = make_static_distributed_tensor(red_dist, dtype=F32)
+             * reduced  = make_static_distributed_tensor(red_dist, dtype=compute_ty)
              * reduced.storage[0] = acc
              * block_tile_reduce_sync(b, reduced, combine, lds_buf=lds, tid, wave_size)
              * total = reduced.storage[0] */
@@ -490,7 +528,7 @@ rocke_kernel_def_t*
             {
                 return NULL;
             }
-            reduced = rocke_make_static_distributed_tensor(b, red_dist, rocke_f32());
+            reduced = rocke_make_static_distributed_tensor(b, red_dist, compute_ty);
             if(reduced == NULL || reduced->num_storage < 1)
             {
                 (void)rocke_i_set_err(b, ROCKE_ERR_VALUE, "reduce2d: empty distributed tensor");
@@ -512,18 +550,23 @@ rocke_kernel_def_t*
             total = rocke_block_lds_reduce(b, acc, lds, tid, BS, combine);
         }
 
-        /* if spec.op == "mean": total = total * rcp(const_f32(float(N))) */
+        /* if spec.op == "mean": total = total * rcp(const_c(float(N))) */
         if(strcmp(spec->op, "mean") == 0)
         {
-            total = rocke_b_fmul(b, total, rocke_b_rcp(b, rocke_b_const_f32(b, (double)N)));
+            total = rocke_b_fmul(
+                b, total, rocke_b_rcp(b, rocke_reduce2d_const_c(b, is_f64, (double)N)));
         }
 
         /* with b.scf_if(b.cmp_eq(tid, const_i32(0))):
-         *     store_scalar_from_f32(b, Y, row, total, dtype=spec.dtype) */
+         *     if is_f64: b.global_store(Y, row, total)
+         *     else: store_scalar_from_f32(b, Y, row, total, dtype=spec.dtype) */
         {
             rocke_if_t iff = rocke_b_scf_if(b, rocke_b_cmp_eq(b, tid, rocke_b_const_i32(b, 0)));
             rocke_b_region_enter(b, iff.then_region);
-            rocke_b_store_scalar_from_f32(b, Y, row, total, spec->dtype);
+            if(is_f64)
+                rocke_b_global_store(b, Y, row, total, 0);
+            else
+                rocke_b_store_scalar_from_f32(b, Y, row, total, spec->dtype);
             rocke_b_region_leave(b);
         }
 
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/ir_model.md b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/ir_model.md
index 4a3f7a87913e..f6ae39dc9af7 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/ir_model.md
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/ir_model.md
@@ -10,7 +10,7 @@ Named scalar types are singletons exported from `core/ir.py`:
 
 ```text
 I1, I8, I32, I64
-F16, BF16, F32
+F16, BF16, F32, F64
 FP8E4M3      # e4m3 fp8 (cvt_fp8_to_f32 / cvt_f32_to_fp8)
 BF8E5M2      # e5m2 fp8 (cvt_bf8_to_f32 / cvt_f32_to_bf8)
 ```
@@ -88,6 +88,7 @@ Kernel attributes consumed by lowering:
 const_i32(value)
 const_i64(value)
 const_f32(value)
+const_f64(value)           # printed as the exact IEEE-754 bit pattern
 fp16_zero()
 zero_vec_f32(n)            # <n x f32>, all zeros
 zero_vec_f32_4()           # <4 x f32>, all zeros
@@ -99,7 +100,7 @@ zero_vec(elem, n)          # f32 / f16 / bf16
 ```text
 add, sub, mul, div, mod                    # integer arithmetic
 fadd, fsub, fmul, fdiv, fneg               # f32 arithmetic
-fmax, fmin                                 # llvm.maxnum / llvm.minnum (f32/f16/bf16)
+fmax, fmin                                 # llvm.maxnum / llvm.minnum (f32/f16/bf16/f64)
 fcmp(pred, a, b)                           # pred in {olt,ole,ogt,oge,oeq,one,ord,uno}
 cmp_lt, cmp_le, cmp_gt, cmp_ge, cmp_eq, cmp_ne   # integer compare (-> i1)
 land, lor, lnot                            # bitwise + i1 logic
@@ -125,7 +126,7 @@ cvt_f32_to_i8_sat(v)     # round + saturate to i8
 
 ### Math intrinsics
 
-LLVM target intrinsics; lowering is f32-centric.
+LLVM target intrinsics; lowering is f32-centric. `exp2`, `sqrt`, `rsqrt` and `tanh` accept only f32 and reject f64 inputs; `rcp` is a plain `fdiv` and accepts f64.
 
 ```text
 exp2(v)        # llvm.exp2.f32
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/optimization/optimization_runbook.md b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/optimization/optimization_runbook.md
index 15ae339a07d8..299a685cbccc 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/optimization/optimization_runbook.md
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/optimization/optimization_runbook.md
@@ -230,7 +230,7 @@ Decide what the kernel computes before deciding how it should run.
   fields explicitly; e.g. `UnifiedAttention2DTiledSpec.dtype="bf16"`,
   `kv_storage_dtype="fp8e4m3"` for the FP8 KV cache path.
 - DSL dtype catalog: `I1`, `I8`, `I16`, `I32`, `I64`, `BF16`, `F16`,
-  `F32`, `FP8E4M3`, `BF8E5M2` (see `core/ir.py:34-97`).
+  `F32`, `F64`, `FP8E4M3`, `BF8E5M2` (see `core/ir.py:34-97`).
 - Quantization helpers in `helpers/quant.py` (`quant_max_abs`,
   `quantize_scalar_f32`, `dequantize_scalar_to_f32`, `quant_ir_type`,
   `ir_to_qdtype`). `QDType` is `Literal["i8", "fp8e4m3", "bf8e5m2"]`.
@@ -242,6 +242,9 @@ Decide what the kernel computes before deciding how it should run.
   norm `<= 5e-3` (was 2.5e-3 before noise widening, see
   `notes/PROPOSALS_IMPLEMENTATION_REPORT.md::F2`), reduce `<= 1.5e-3`,
   gemm `<= 7e-2`).
+- fp64 kernels compute natively in f64 and are checked against a float64
+  reference (`tests/instances/differential/numeric.py`: elementwise
+  `<= 1e-12`, reduce `<= 1e-10`).
 - For fp16 inputs with fp32 accumulation over O(100) terms, expect
   errors near the fp16 ULP floor for correct kernels. Errors two
   orders of magnitude higher almost always indicate structural bugs
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/primitives/wave_and_cross_lane.md b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/primitives/wave_and_cross_lane.md
index bd7b325e2173..424c82cdb8b9 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/primitives/wave_and_cross_lane.md
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/primitives/wave_and_cross_lane.md
@@ -87,7 +87,7 @@ Notes:
 b.warp_shuffle_xor(v, lane_xor)
 ```
 
-Convenience wrapper: lane `l` gets `v` from lane `l ^ lane_xor`. Works for f32 and i32 directly; bitcast first for half / bf16. Internally:
+Convenience wrapper: lane `l` gets `v` from lane `l ^ lane_xor`. Works for f32 and i32 directly; bitcast first for half / bf16. An f64 `v` is split into two i32 halves, each shuffled separately, and reassembled. Internally:
 
 ```text
 addr = (lane_id() ^ lane_xor) << 2
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md
index 4730a9f37f04..3341a2835204 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md
@@ -6,7 +6,7 @@ Complete reference of operations recognized by `core/ir.py` and lowered to AMDGP
 
 | Op name                | Builder                           | LLVM emission                          |
 |------------------------|-----------------------------------|----------------------------------------|
-| `arith.constant`       | `const_i32`, `const_i64`, `const_f32`, `fp16_zero` | constant value |
+| `arith.constant`       | `const_i32`, `const_i64`, `const_f32`, `const_f64`, `fp16_zero` | constant value |
 | `arith.constant_vec`   | `zero_vec_f32`, `zero_vec`        | `zeroinitializer` vector               |
 | `arith.add`            | `add`                             | `add` (integer)                        |
 | `arith.sub`            | `sub`                             | `sub`                                  |
@@ -51,6 +51,8 @@ Complete reference of operations recognized by `core/ir.py` and lowered to AMDGP
 | `math.tanh`    | `tanh`       | piecewise f32 expansion: OCML polynomial for `abs(x) < 0.625`, otherwise `exp2` + reciprocal; bitwise sign restoration |
 | `math.rcp`     | `rcp`        | `1.0 / v` (hardware reciprocal) |
 
+`exp2`, `sqrt`, `rsqrt` and `tanh` are f32-only and reject f64 inputs; `rcp` accepts f64.
+
 `clamp_f32(v, lo, hi)` is `fmin(hi, fmax(lo, v))` — folds to `v_med3_f32`.
 
 ## Vector
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/__init__.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/__init__.py
index 2f92e30dee83..266997b8eb4b 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/__init__.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/__init__.py
@@ -62,6 +62,7 @@
     BF16,
     F16,
     F32,
+    F64,
     FP8E4M3,
     I1,
     I8,
@@ -168,6 +169,7 @@
     "BF16",
     "F16",
     "F32",
+    "F64",
     "FP8E4M3",
     "I1",
     "I8",
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/__init__.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/__init__.py
index b875e9036502..59c27cc93da1 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/__init__.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/__init__.py
@@ -52,6 +52,7 @@
     BF16,
     F16,
     F32,
+    F64,
     FP8E4M3,
     FP4E2M1,
     FP6E2M3,
@@ -121,6 +122,7 @@
     "BF16",
     "F16",
     "F32",
+    "F64",
     "FP8E4M3",
     "I1",
     "I8",
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/dtypes.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/dtypes.py
index 731eafe5efa6..86df9d5d5ef0 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/dtypes.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/dtypes.py
@@ -25,6 +25,9 @@
     "fp32": "fp32",
     "tf32": "tf32",
     "xf32": "tf32",
+    "f64": "fp64",
+    "double": "fp64",
+    "fp64": "fp64",
     "fp8": "fp8e4m3",
     "fp8e4m3": "fp8e4m3",
     "e4m3": "fp8e4m3",
@@ -77,6 +80,7 @@ class DTypeInfo:
         (
             DTypeCategory.FLOAT,
             (
+                ("fp64", 64),
                 ("fp32", 32),
                 ("tf32", 32),
                 ("fp16", 16),
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/ir.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/ir.py
index 4880390aa68d..e5d4e330286e 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/ir.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/ir.py
@@ -53,6 +53,7 @@ def __repr__(self) -> str:
 F16 = Type("f16")
 F32 = Type("f32")
 TF32 = Type("tf32")
+F64 = Type("f64")
 FP8E4M3 = Type("fp8e4m3")
 BF8E5M2 = Type("bf8e5m2")
 FP4E2M1 = Type("fp4e2m1")
@@ -78,6 +79,7 @@ def dtype_to_ir_type(dtype: str) -> Type:
         "bf16": BF16,
         "fp32": F32,
         "tf32": TF32,
+        "fp64": F64,
         "fp8e4m3": FP8E4M3,
         "bf8e5m2": BF8E5M2,
         "fp4e2m1": FP4E2M1,
@@ -750,6 +752,15 @@ def const_f32(self, value: float) -> Value:
         )
         return op.result
 
+    def const_f64(self, value: float) -> Value:
+        op = self._op(
+            "arith.constant",
+            result_types=[F64],
+            attrs={"value": float(value), "ity": "f64"},
+            result_name_hint="c",
+        )
+        return op.result
+
     def add(self, a: Value, b: Value) -> Value:
         return self._op("arith.add", [a, b], [a.type], result_name_hint="add").result
 
@@ -1776,9 +1787,10 @@ def global_load_vN(
         """Vectorised global load of N consecutive values.
 
         Supports f16/bf16/i16 (N in {2, 4, 6, 8, 16}), f32/i32/tf32
-        (N in {2, 3, 4, 8}), and fp8e4m3/bf8e5m2/i8 (N in {2, 4, 8, 12, 16}).
-        Loads exactly N elements. Instruction selection depends on target and
-        alignment; 96-bit payloads do not require a 96-bit scalar type.
+        (N in {2, 3, 4, 8}), fp8e4m3/bf8e5m2/i8 (N in {2, 4, 8, 12, 16}), and
+        f64 (N = 2). Loads exactly N elements. Instruction selection depends on
+        target and alignment; 96-bit payloads do not require a 96-bit scalar
+        type.
 
         Default alignment is the payload size for power-of-two loads, and
         element alignment for 12-byte loads. An explicit alignment is a caller
@@ -1813,9 +1825,15 @@ def global_load_vN(
                 raise ValueError(
                     f"unsupported vector width for {dtype.name} global_load_vN: {n}"
                 )
+        elif dtype.name == "f64":
+            elem_bytes = 8
+            if n != 2:
+                raise ValueError(
+                    f"unsupported vector width for {dtype.name} global_load_vN: {n}"
+                )
         else:
             raise ValueError(
-                "global_load_vN supports f16/bf16/i16/f32/i32/tf32/fp8e4m3/bf8e5m2/i8, "
+                "global_load_vN supports f16/bf16/i16/f32/i32/tf32/fp8e4m3/bf8e5m2/i8/f64, "
                 f"got {dtype.name}"
             )
         attrs = {
@@ -1992,19 +2010,20 @@ def smem_store_vN(
         if not isinstance(value.type, VectorType):
             raise ValueError("smem_store_vN expects vector value for n > 1")
         elem_name = value.type.elem.name
-        allowed_n = (
-            (2, 4, 8, 16) if elem_name in ("i8", "fp8e4m3", "bf8e5m2") else (2, 4, 8)
-        )
+        if elem_name in ("i8", "fp8e4m3", "bf8e5m2"):
+            allowed_n = (2, 4, 8, 16)
+            elem_bytes = 1
+        elif elem_name == "f64":
+            allowed_n = (2,)
+            elem_bytes = 8
+        else:
+            allowed_n = (2, 4, 8)
+            elem_bytes = 4 if elem_name in ("f32", "i32", "tf32") else 2
         if n not in allowed_n:
             raise ValueError(
                 f"unsupported vector width for smem_store_vN of {elem_name}: {n} "
                 f"(allowed: {allowed_n})"
             )
-        elem_bytes = (
-            1
-            if elem_name in ("i8", "fp8e4m3", "bf8e5m2")
-            else 4 if elem_name in ("f32", "i32", "tf32") else 2
-        )
         self._op(
             "tile.smem_store_vN",
             [smem, *indices, value],
@@ -2036,10 +2055,11 @@ def smem_load_vN_f16(self, smem: Value, *indices, n: int = 0) -> Value:
 
     def smem_load_vN(self, smem: Value, *indices, dtype: Type, n: int = 0) -> Value:
         """LDS load of ``<N x dtype>``. Supports 8-bit (fp8e4m3 / bf8e5m2 /
-        i8), 16-bit (f16 / bf16) and 32-bit (f32 / i32) element types;
-        loads exactly N elements. In addition to power-of-two widths, accepts
-        96-bit payloads (12 bytes, six halfwords, or three words), using element
-        alignment. The target and alignment determine instruction selection.
+        i8), 16-bit (f16 / bf16), 32-bit (f32 / i32) and 64-bit (f64, ``n in
+        {1, 2}``) element types; loads exactly N elements. In addition to
+        power-of-two widths, accepts 96-bit payloads (12 bytes, six halfwords,
+        or three words), using element alignment. The target and alignment
+        determine instruction selection.
         """
         if dtype.name not in (
             "f16",
@@ -2050,16 +2070,20 @@ def smem_load_vN(self, smem: Value, *indices, dtype: Type, n: int = 0) -> Value:
             "fp8e4m3",
             "bf8e5m2",
             "i8",
+            "f64",
         ):
             raise ValueError(
                 "smem_load_vN supports f16 / bf16 / f32 / i32 / tf32 / fp8e4m3 / "
-                f"bf8e5m2 / i8, got {dtype.name}"
+                f"bf8e5m2 / i8 / f64, got {dtype.name}"
             )
-        allowed_n = (
-            (1, 2, 4, 8, 12, 16)
-            if dtype.name in ("fp8e4m3", "bf8e5m2", "i8")
-            else (1, 2, 4, 6, 8) if dtype.name in ("f16", "bf16") else (1, 2, 3, 4, 8)
-        )
+        if dtype.name in ("fp8e4m3", "bf8e5m2", "i8"):
+            allowed_n = (1, 2, 4, 8, 12, 16)
+        elif dtype.name in ("f16", "bf16"):
+            allowed_n = (1, 2, 4, 6, 8)
+        elif dtype.name == "f64":
+            allowed_n = (1, 2)
+        else:
+            allowed_n = (1, 2, 3, 4, 8)
         if n not in allowed_n:
             raise ValueError(
                 f"unsupported vector width {n} for smem_load_vN of {dtype.name} "
@@ -3336,8 +3360,14 @@ def warp_shuffle_xor(self, v: Value, lane_xor: int) -> Value:
         permlane32_swap or ds_bpermute).
 
         Works for any 32-bit scalar `v` (f32, i32). For half/bfloat,
-        bitcast to i32 via a 2-element vector first.
+        bitcast to i32 via a 2-element vector first. An f64 `v` is split
+        into two i32 halves, each shuffled on the 32-bit path.
         """
+        if v.type.name == "f64":
+            halves = self.bitcast(v, VectorType(I32, 2))
+            lo = self.warp_shuffle_xor(self.vec_extract(halves, 0), lane_xor)
+            hi = self.warp_shuffle_xor(self.vec_extract(halves, 1), lane_xor)
+            return self.bitcast(self.vec_pack([lo, hi], I32), F64)
         if 1 <= lane_xor <= 31:
             # Emits `ds_swizzle` (XOR pattern in the immediate offset), NOT
             # `ds_bpermute`: 1 LDS op, no addr-compute. This is the path the
@@ -4441,7 +4471,8 @@ def global_store_vN(
 
         Supports the full element-type catalog the LLVM lowering already
         emits: ``f16`` / ``bf16`` / ``i16`` (2-byte), ``f32`` / ``i32`` / ``tf32``
-        (4-byte), ``i8`` / ``fp8e4m3`` / ``bf8e5m2`` (1-byte). Lowers to
+        (4-byte), ``i8`` / ``fp8e4m3`` / ``bf8e5m2`` (1-byte), ``f64``
+        (8-byte, N in {1, 2}). Lowers to
         a single ``store <N x elem>`` with the supplied address alignment.
         Payload width and address alignment are independent; target and
         alignment determine whether the transfer uses one machine instruction.
@@ -4469,9 +4500,13 @@ def global_store_vN(
                 raise ValueError(f"global_store_vN n=16 not supported for {elem_name}")
         elif elem_name in ("i8", "fp8e4m3", "bf8e5m2"):
             elem_bytes = 1
+        elif elem_name == "f64":
+            elem_bytes = 8
+            if n > 2:
+                raise ValueError(f"global_store_vN n={n} not supported for {elem_name}")
         else:
             raise ValueError(
-                "global_store_vN supports f16/bf16/i16/f32/i32/tf32/i8/fp8e4m3/bf8e5m2, "
+                "global_store_vN supports f16/bf16/i16/f32/i32/tf32/i8/fp8e4m3/bf8e5m2/f64, "
                 f"got {elem_name}"
             )
         attrs = {
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py
index a18b3d9a6fcb..a3a3a10028cb 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py
@@ -46,6 +46,7 @@
     "f16": "fp16",
     "bf16": "bf16",
     "f32": "float",
+    "f64": "double",
     "fp8e4m3": "fp8e4m3",
     "bf8e5m2": "bf8e5m2",
 }
@@ -73,6 +74,8 @@ def _type_to_hip(t) -> str:
             return f"bf16x{t.count}"
         if elem == "f32":
             return f"f32x{t.count}"
+        if elem == "f64":
+            return f"f64x{t.count}"
         if elem in ("i32", "tf32"):
             return f"i32x{t.count}"
         if elem == "i16":
@@ -128,6 +131,7 @@ def _type_to_hip(t) -> str:
 _ROCKE_VEC(bf16, bf16x, 8); _ROCKE_VEC(bf16, bf16x, 16);
 _ROCKE_VEC(float, f32x, 1); _ROCKE_VEC(float, f32x, 2); _ROCKE_VEC(float, f32x, 4);
 _ROCKE_VEC(float, f32x, 8); _ROCKE_VEC(float, f32x, 16);
+_ROCKE_VEC(double, f64x, 1); _ROCKE_VEC(double, f64x, 2);
 _ROCKE_VEC(int, i32x, 1); _ROCKE_VEC(int, i32x, 2); _ROCKE_VEC(int, i32x, 3);
 _ROCKE_VEC(int, i32x, 4); _ROCKE_VEC(int, i32x, 8); _ROCKE_VEC(int, i32x, 16);
 _ROCKE_VEC(int16_t, i16x, 1); _ROCKE_VEC(int16_t, i16x, 2);
@@ -187,6 +191,7 @@ def _name(v: Value) -> str:
     "f16": "f16x",
     "bf16": "bf16x",
     "f32": "f32x",
+    "f64": "f64x",
     "i32": "i32x",
     "tf32": "i32x",
     "i16": "i16x",
@@ -234,6 +239,22 @@ def _f32_literal(val: float) -> str:
     return f"{val}f"
 
 
+def _f64_literal(val: float) -> str:
+    """Format a Python float for C++ double literal context.
+
+    ``repr`` is the shortest round-trip form, so finite values are exact;
+    ``inf`` / ``-inf`` / ``nan`` use the ``<math.h>`` macros as in
+    :func:`_f32_literal`.
+    """
+    import math
+
+    if math.isnan(val):
+        return "((double)NAN)"
+    if math.isinf(val):
+        return "((double)-INFINITY)" if val < 0 else "((double)INFINITY)"
+    return repr(val)
+
+
 def _encode_waitcnt_gfx9_10(vmcnt: int, expcnt: int, lgkmcnt: int) -> int:
     """Encode `s_waitcnt` for the gfx9/gfx10 layout used by gfx950.
 
@@ -372,6 +393,8 @@ def _op_arith_constant(self, op: Op) -> None:
                 self._emit(f"{cpp_t} {_name(res)} = (fp16){literal};")
             else:
                 self._emit(f"{cpp_t} {_name(res)} = {literal};")
+        elif ity == "f64":
+            self._emit(f"{cpp_t} {_name(res)} = {_f64_literal(float(val))};")
         else:
             self._emit(f"{cpp_t} {_name(res)} = {val};")
 
@@ -1440,11 +1463,14 @@ def _op_math_log2(self, op: Op) -> None:
 
     def _op_math_rcp(self, op: Op) -> None:
         # AMDGPU has a hardware reciprocal; emit the builtin directly for
-        # f32, promote-compute-demote for f16/bf16.
+        # f32, promote-compute-demote for f16/bf16. f64 divides in double
+        # (matching the LLVM ``fdiv``) so the result keeps f64 precision.
         (v,) = op.operands
         tname = op.result.type.name
         cpp_t = _type_to_hip(op.result.type)
-        if tname == "f32":
+        if tname == "f64":
+            self._emit(f"{cpp_t} {_name(op.result)} = 1.0 / {_name(v)};")
+        elif tname == "f32":
             self._emit(
                 f"{cpp_t} {_name(op.result)} = __builtin_amdgcn_rcpf({_name(v)});"
             )
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py
index 34dac431b950..7d1660ea3eaf 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py
@@ -419,9 +419,11 @@ def _resolve_llvm_flavor() -> str:
     "maxnum.f32": "declare float @llvm.maxnum.f32(float, float)",
     "maxnum.f16": "declare half @llvm.maxnum.f16(half, half)",
     "maxnum.bf16": "declare bfloat @llvm.maxnum.bf16(bfloat, bfloat)",
+    "maxnum.f64": "declare double @llvm.maxnum.f64(double, double)",
     "minnum.f32": "declare float @llvm.minnum.f32(float, float)",
     "minnum.f16": "declare half @llvm.minnum.f16(half, half)",
     "minnum.bf16": "declare bfloat @llvm.minnum.bf16(bfloat, bfloat)",
+    "minnum.f64": "declare double @llvm.minnum.f64(double, double)",
     "fabs.f32": "declare float @llvm.fabs.f32(float)",
     "fabs.f16": "declare half @llvm.fabs.f16(half)",
     "fabs.bf16": "declare bfloat @llvm.fabs.bf16(bfloat)",
@@ -1048,6 +1050,8 @@ def _llvm_type(t: Type) -> str:
         return "i8"
     if t.name == "f32":
         return "float"
+    if t.name == "f64":
+        return "double"
     raise NotImplementedError(f"no LLVM mapping for type {t!r}")
 
 
@@ -1565,6 +1569,8 @@ def _operand(self, v: Value) -> str:
                 return _fp32_hex(val)
             if ity == "f16":
                 return _fp16_hex(val)
+            if ity == "f64":
+                return _fp64_hex(val)
             return str(int(val))
         return v.name
 
@@ -1852,6 +1858,7 @@ def _compute_smem_layout(self) -> None:
             "tf32": 4,
             "f32": 4,
             "i64": 8,
+            "f64": 8,
         }
 
         pool_name = f"@smem_pool.{self.kernel.name}"
@@ -1870,6 +1877,8 @@ def _seg_size(stype: "SmemType") -> int:
             return seg
 
         def _align(stype: "SmemType") -> int:
+            if stype.elem.name == "f64":
+                return 8
             return 16 if stype.elem.name in ("i8", "fp8e4m3", "bf8e5m2") else 4
 
         # ---- live intervals from the kernel body ----
@@ -2238,11 +2247,17 @@ def _op_arith_fmax(self, op: Op) -> None:
         # Dispatch on the operand FP width: half, float, bfloat each
         # have their own maxnum intrinsic. The previous f32-only path
         # silently mis-typed half operands as float and broke comgr.
-        llvm_ty = {"f32": "float", "f16": "half", "bf16": "bfloat"}.get(ty_name)
+        llvm_ty = {
+            "f32": "float",
+            "f16": "half",
+            "bf16": "bfloat",
+            "f64": "double",
+        }.get(ty_name)
         intrin_key = {
             "f32": "maxnum.f32",
             "f16": "maxnum.f16",
             "bf16": "maxnum.bf16",
+            "f64": "maxnum.f64",
         }.get(ty_name)
         if llvm_ty is None or intrin_key is None:
             raise NotImplementedError(f"fmax: unsupported FP type {ty_name!r}")
@@ -2255,11 +2270,17 @@ def _op_arith_fmax(self, op: Op) -> None:
     def _op_arith_fmin(self, op: Op) -> None:
         a, b = op.operands
         ty_name = a.type.name
-        llvm_ty = {"f32": "float", "f16": "half", "bf16": "bfloat"}.get(ty_name)
+        llvm_ty = {
+            "f32": "float",
+            "f16": "half",
+            "bf16": "bfloat",
+            "f64": "double",
+        }.get(ty_name)
         intrin_key = {
             "f32": "minnum.f32",
             "f16": "minnum.f16",
             "bf16": "minnum.bf16",
+            "f64": "minnum.f64",
         }.get(ty_name)
         if llvm_ty is None or intrin_key is None:
             raise NotImplementedError(f"fmin: unsupported FP type {ty_name!r}")
@@ -3114,7 +3135,7 @@ def _op_tile_smem_store(self, op: Op) -> None:
             f"{', '.join(gidx)}"
         )
         # Alignment is the element byte size: 1 for i8, 2 for f16/bf16,
-        # 4 for f32/i32, 8 for i64. The AMDGPU backend rejects loads /
+        # 4 for f32/i32, 8 for i64/f64. The AMDGPU backend rejects loads /
         # stores with under-aligned addresses on ``addrspace(3)``.
         align = {
             "i8": 1,
@@ -3126,6 +3147,7 @@ def _op_tile_smem_store(self, op: Op) -> None:
             "tf32": 4,
             "f32": 4,
             "i64": 8,
+            "f64": 8,
         }.get(value.type.name, 2)
         self._current().emit(
             f"  store {_llvm_type(value.type)} {self._operand(value)}, ptr addrspace(3) {gep}, align {align}"
@@ -3191,6 +3213,7 @@ def _op_tile_smem_store_vN(self, op: Op) -> None:
             "tf32": 4,
             "f32": 4,
             "i64": 8,
+            "f64": 8,
         }.get(
             value.type.elem.name, 2
         )  # type: ignore[attr-defined]
@@ -3265,7 +3288,8 @@ def _op_tile_smem_load_vN(self, op: Op) -> None:
         )
         elem_ty = _llvm_type(op.result.type.elem)  # type: ignore[attr-defined]
         # Element byte size drives the vector alignment. 16-bit
-        # (f16 / bf16): 2 bytes; 32-bit (f32 / i32): 4 bytes.
+        # (f16 / bf16): 2 bytes; 32-bit (f32 / i32): 4 bytes; 64-bit
+        # (i64 / f64): 8 bytes.
         elem_bytes = {
             "i8": 1,
             "f16": 2,
@@ -3274,6 +3298,7 @@ def _op_tile_smem_load_vN(self, op: Op) -> None:
             "tf32": 4,
             "f32": 4,
             "i64": 8,
+            "f64": 8,
         }.get(
             op.result.type.elem.name,
             2,  # type: ignore[attr-defined]
@@ -5509,6 +5534,7 @@ def _op_memref_global_store_vN(self, op: Op) -> None:
             "tf32": 4,
             "f32": 4,
             "i64": 8,
+            "f64": 8,
         }.get(elem_name, 2)
         align = int(op.attrs.get("align", vec * elem_bytes))
         if align <= 0 or align & (align - 1):
@@ -6407,6 +6433,14 @@ def _fp16_hex(x: float) -> str:
     return f"0xH{bits:04X}"
 
 
+def _fp64_hex(x: float) -> str:
+    # LLVM textual IR spells `double` hex constants as the raw IEEE-754 bits.
+    import struct
+
+    bits = struct.unpack("<Q", struct.pack("<d", float(x)))[0]
+    return f"0x{bits:016X}"
+
+
 def _lower_kernel_to_llvm_python(
     kernel: KernelDef,
     *,
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/distribution.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/distribution.py
index eb3fcc7ecd97..da07477d774d 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/distribution.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/distribution.py
@@ -46,7 +46,7 @@
 from itertools import product
 from typing import Callable, Iterable, List, Optional, Sequence, Tuple
 
-from ..core.ir import F32, IRBuilder, Type, Value
+from ..core.ir import F32, F64, IRBuilder, Type, Value
 from .tensor_view import TileWindow
 
 
@@ -894,6 +894,7 @@ def load_tile(
     works as well as f16/bf16: the f32 promote is a no-op for f32 and
     the Y slots simply hold the native f32 scalars. This lets the
     reduce / norm f32-accumulator tiles flow through the same path.
+    ``f64`` computes natively, so its Y slots hold unpromoted f64 scalars.
 
     **Masked / OOB-safe variant** (``mask_fn``): when ``N`` need not divide
     the block extent (img2col K-tail, pooling, topk) the in-bounds region
@@ -908,9 +909,9 @@ def load_tile(
     """
     if traits is None:
         traits = make_load_store_traits(distribution)
-    if window.dtype.name not in ("f16", "bf16", "f32"):
+    if window.dtype.name not in ("f16", "bf16", "f32", "f64"):
         raise NotImplementedError(
-            f"load_tile dtype {window.dtype.name} not wired (f16/bf16/f32 only)"
+            f"load_tile dtype {window.dtype.name} not wired (f16/bf16/f32/f64 only)"
         )
     if mask_fn is not None and window.view.addr_space != "buffer":
         raise NotImplementedError(
@@ -941,6 +942,8 @@ def load_tile(
             scalars = [
                 v if window.dtype.name == "f32" else b.cast_to_f32(v) for v in raw
             ]
+        elif window.dtype.name == "f64":
+            scalars = window.load_vec_scalars(b, *x_coords, n=traits.scalar_per_vector)
         else:
             scalars = window.load_vec_as_f32(b, *x_coords, n=traits.scalar_per_vector)
         # Splice each loaded scalar into the matching Y slot along
@@ -966,20 +969,20 @@ def store_tile(
     ``distributed``, packs them into ``traits.scalar_per_vector``-wide
     vectors, and writes them through ``window``.
 
-    Like :func:`load_tile`, ``f32`` is supported in addition to
-    f16/bf16. For f16/bf16 we route through
-    :meth:`TileWindow.store_vec_from_f32` (f32 demote + pack); for f32
-    the demote is a no-op so we pack the native scalars directly via
+    Like :func:`load_tile`, ``f32`` and ``f64`` are supported in addition
+    to f16/bf16. For f16/bf16 we route through
+    :meth:`TileWindow.store_vec_from_f32` (f32 demote + pack); for f32 /
+    f64 there is no demote so we pack the native scalars directly via
     :meth:`TileWindow.store_vec` / :meth:`TileWindow.store_scalar`.
     """
     distribution = distributed.distribution
     if traits is None:
         traits = make_load_store_traits(distribution)
     quant_dtypes = ("i8", "fp8e4m3", "bf8e5m2")
-    if window.dtype.name not in ("f16", "bf16", "f32", "i32") + quant_dtypes:
+    if window.dtype.name not in ("f16", "bf16", "f32", "f64", "i32") + quant_dtypes:
         raise NotImplementedError(
             f"store_tile dtype {window.dtype.name} not wired "
-            "(f16/bf16/f32/i32/i8/fp8e4m3/bf8e5m2)"
+            "(f16/bf16/f32/f64/i32/i8/fp8e4m3/bf8e5m2)"
         )
     for y_base in traits.iterate_accesses():
         x_coords = distribution.calculate_x(
@@ -1008,10 +1011,10 @@ def store_tile(
             else:
                 packed = b.vec_pack(scalars, window.dtype)
                 window.store_vec(b, *x_coords, value=packed, n=len(scalars))
-        elif window.dtype.name == "f32":
-            # f32 demote is a no-op; TileWindow.store_vec_from_f32
-            # rejects non-16-bit dtypes, so pack/store the native f32
-            # scalars directly here.
+        elif window.dtype.name in ("f32", "f64"):
+            # f32 demote is a no-op and f64 computes natively;
+            # TileWindow.store_vec_from_f32 rejects non-16-bit dtypes, so
+            # pack/store the native scalars directly here.
             if len(scalars) == 1:
                 window.store_scalar(b, *x_coords, value=scalars[0])
             else:
@@ -1493,7 +1496,10 @@ def _fold(x: Value, y: Value) -> Value:
         for i in range(thread_buf_size):
             v = reduced.storage[i]
             slot = b.add(warp, b.const_i32(i * num_reduce_warps))
-            b.smem_store_vN_f32(lds_buf, [slot], v, 1)
+            if v.type.name == "f64":
+                b.smem_store_vN(lds_buf, [slot], v, 1)
+            else:
+                b.smem_store_vN_f32(lds_buf, [slot], v, 1)
     b.sync()
     # local_warp_id groups num_reduce_warps consecutive warps; each reads
     # its group's partials and tree-folds them.
@@ -1503,7 +1509,10 @@ def _fold(x: Value, y: Value) -> Value:
         parts: List[Value] = []
         for idx in range(num_reduce_warps):
             slot = b.add(local_smem_os, b.const_i32(i * num_reduce_warps + idx))
-            vec = b.smem_load_vN_f32(lds_buf, slot, n=1)
+            if reduced.storage[i].type.name == "f64":
+                vec = b.smem_load_vN(lds_buf, slot, dtype=F64, n=1)
+            else:
+                vec = b.smem_load_vN_f32(lds_buf, slot, n=1)
             parts.append(b.vec_extract(vec, 0))
         # Pairwise power-of-two tree fold (matches the C++ stride doubling).
         while len(parts) > 1:
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/io.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/io.py
index 379484c5ab18..3ace81fa7e94 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/io.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/io.py
@@ -22,7 +22,7 @@
 
 from typing import Literal
 
-from ..core.ir import BF16, F16, IRBuilder, TemporalHint, Type, Value
+from ..core.ir import BF16, F16, F64, IRBuilder, TemporalHint, Type, Value
 
 
 __all__ = [
@@ -51,15 +51,17 @@
 def io_ir_type(dtype: str) -> Type:
     """Map a dtype string to the canonical IR type object.
 
-    Accepts ``"f16"``, ``"fp16"`` (alias), or ``"bf16"``. Raises
-    :class:`ValueError` for any other value -- f8 and i8 paths go
+    Accepts ``"f16"``, ``"fp16"`` (alias), ``"bf16"``, or ``"f64"``.
+    Raises :class:`ValueError` for any other value -- f8 and i8 paths go
     through their own helpers because their compute dtype isn't f32.
     """
     if dtype in ("f16", "fp16"):
         return F16
     if dtype == "bf16":
         return BF16
-    raise ValueError(f"unsupported I/O dtype {dtype!r}; expected f16/fp16/bf16")
+    if dtype == "f64":
+        return F64
+    raise ValueError(f"unsupported I/O dtype {dtype!r}; expected f16/fp16/bf16/f64")
 
 
 def load_scalar(b: IRBuilder, ptr: Value, idx: Value, *, dtype: str) -> Value:
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/reduction.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/reduction.py
index b36bf3c96f58..5eff1be0df22 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/reduction.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/reduction.py
@@ -29,7 +29,7 @@
 
 from typing import Callable, List, Literal, Tuple
 
-from ..core.ir import F32, I32, IRBuilder, Value
+from ..core.ir import F32, F64, I32, IRBuilder, Value
 
 
 __all__ = [
@@ -316,15 +316,16 @@ def block_lds_reduce_with_wave_prologue(
     on small-shape reductions in the working ``instances/reduce.py``
     prototype.
 
-    ``lds_buf`` must point to at least ``num_warps`` f32 slots; reuse
-    the kernel's existing ``block_size``-element LDS allocation so
-    the kernel's LDS footprint doesn't change.
+    ``lds_buf`` must point to at least ``num_warps`` slots of ``val``'s
+    type (f32 or f64); reuse the kernel's existing ``block_size``-element
+    LDS allocation so the kernel's LDS footprint doesn't change.
 
     Promoted from ``instances/reduce.py::_block_tile_reduce`` (P20).
     """
-    if val.type.name != "f32":
+    if val.type.name not in ("f32", "f64"):
         raise ValueError(
-            f"block_lds_reduce_with_wave_prologue expects f32 input, got {val.type.name}"
+            "block_lds_reduce_with_wave_prologue expects f32 or f64 input, "
+            f"got {val.type.name}"
         )
 
     warp_partial = _warp_xor_reduce(b, val, combine=combine, wave_size=wave_size)
@@ -336,13 +337,20 @@ def block_lds_reduce_with_wave_prologue(
     c_wave = b.const_i32(wave_size)
     lane = b.mod(tid, c_wave)
     warp = b.div(tid, c_wave)
+    is_f64 = val.type.name == "f64"
     with b.scf_if(b.cmp_eq(lane, b.const_i32(0))):
-        b.smem_store_vN_f32(lds_buf, [warp], warp_partial, 1)
+        if is_f64:
+            b.smem_store_vN(lds_buf, [warp], warp_partial, 1)
+        else:
+            b.smem_store_vN_f32(lds_buf, [warp], warp_partial, 1)
     b.sync()
 
     parts: List[Value] = []
     for w in range(num_warps):
-        v_vec = b.smem_load_vN_f32(lds_buf, b.const_i32(w), n=1)
+        if is_f64:
+            v_vec = b.smem_load_vN(lds_buf, b.const_i32(w), dtype=F64, n=1)
+        else:
+            v_vec = b.smem_load_vN_f32(lds_buf, b.const_i32(w), n=1)
         parts.append(b.vec_extract(v_vec, 0))
     return _tree_reduce_scalars(b, combine, parts)
 
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/sweep.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/sweep.py
index fafa046b9382..790711b04bd5 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/sweep.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/sweep.py
@@ -95,7 +95,8 @@ def sweep_row_chunks(
 
     The :class:`TileWindow` ``tile`` must be 2D with the row dim first;
     each call to ``tile.load_vec_as_f32(b, 0, n_off, n=vec)`` reads a
-    contiguous vec-wide slice of the current row, promoted to f32. The
+    contiguous vec-wide slice of the current row, promoted to f32 (f64
+    tiles are read without promotion). The
     ``row`` argument shifts the tile's row origin in-place (matching
     CK Tile's ``set_window_origin``); pass ``None`` to leave the
     current origin alone (e.g. when the caller already moved the tile
@@ -140,7 +141,10 @@ def sweep_row_chunks(
             b.mul(b.const_i32(k * block_size), c_vec),
             b.mul(tid, c_vec),
         )
-        x_scalars = tile.load_vec_as_f32(b, b.const_i32(0), n_off, n=vec)
+        if tile.dtype.name == "f64":
+            x_scalars = tile.load_vec_scalars(b, b.const_i32(0), n_off, n=vec)
+        else:
+            x_scalars = tile.load_vec_as_f32(b, b.const_i32(0), n_off, n=vec)
         if cache:
             cached.extend(x_scalars)
         if body is not None:
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/tensor_view.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/tensor_view.py
index 9e33e2f5436c..ba1a92a95d5e 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/tensor_view.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/tensor_view.py
@@ -421,8 +421,8 @@ def store_scalar(
     def load_vec(self, b: IRBuilder, indices: Sequence[Value], n: int) -> Value:
         """Vectorised load of ``n`` consecutive elements starting at
         ``indices``. Supports ``n in {2, 4, 8}`` for f16/bf16 (global &
-        LDS); buffer ops use ``dwords = n // 2`` for f16 and support
-        ``n in {2, 4, 8}`` accordingly."""
+        LDS) and ``n = 2`` for f64 (global); buffer ops use ``dwords =
+        n // 2`` for f16 and support ``n in {2, 4, 8}`` accordingly."""
         if self.addr_space == "lds":
             if self.dtype.name in ("f16", "bf16"):
                 return b.smem_load_vN(self.base, *indices, dtype=self.dtype, n=n)
@@ -447,7 +447,7 @@ def load_vec(self, b: IRBuilder, indices: Sequence[Value], n: int) -> Value:
                 f"buffer vec load not yet wired for dtype {self.dtype.name}"
             )
         off = self.desc.offset(b, indices)
-        if self.dtype.name in ("f16", "bf16"):
+        if self.dtype.name in ("f16", "bf16", "f64"):
             return b.global_load_vN(self.base, off, self.dtype, n)
         if self.dtype.name == "f32":
             # f32 global vec loads aren't wired through ``global_load_vN``
@@ -880,6 +880,20 @@ def store_vec(
     ) -> None:
         self.view.store_vec(b, self._global_indices(b, local_indices), value=value, n=n)
 
+    def load_vec_scalars(
+        self, b: IRBuilder, *local_indices: Value, n: int
+    ) -> list[Value]:
+        """Vector load split into ``n`` scalars in the window's own dtype.
+
+        The no-promotion counterpart of :meth:`load_vec_as_f32`, used when
+        compute stays in the storage dtype (f64). ``n == 1`` routes through
+        :meth:`load_scalar`.
+        """
+        if n == 1:
+            return [self.load_scalar(b, *local_indices)]
+        v = self.load_vec(b, *local_indices, n=n)
+        return [b.vec_extract(v, i) for i in range(n)]
+
     # ---- compute-promoting vector ops ----
 
     def load_vec_as_f32(
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/SUPPORT_MATRIX.md b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/SUPPORT_MATRIX.md
index 9c3377b3c2f6..434ef8a313ad 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/SUPPORT_MATRIX.md
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/SUPPORT_MATRIX.md
@@ -40,6 +40,10 @@ atom (CDNA: `pipeline=mem`/`compv*`; gfx1151: `pipeline=mem`, `epilogue=default`
 
 These emit generic AMDGPU IR; arch only sets the comgr target triple.
 
+`elementwise` and `reduce2d` also accept `dtype="f64"` (native f64
+compute, `vec=2`). f64 elementwise is limited to the exact ops `copy`,
+`neg`, `abs`, `relu`, `add`, `sub`, `mul`, `max` and `min`.
+
 ---
 
 ## gfx1250 native scaled GEMM
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/common/elementwise.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/common/elementwise.py
index 0b8d0e1598b8..17350724c7fb 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/common/elementwise.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/common/elementwise.py
@@ -13,7 +13,8 @@
 * Unary ops: ``copy``, ``neg``, ``abs``, ``relu``, ``gelu_tanh``, ``silu``, ``exp2``,
   ``tanh``
 * Binary ops: ``add``, ``sub``, ``mul``, ``max``, ``min``
-* Dtypes: ``f16`` and ``bf16`` for I/O (compute is f32 internally)
+* Dtypes: ``f16`` and ``bf16`` for I/O (compute is f32 internally), and
+  ``f64`` (computed natively; exact ops only, ``vec=2``)
 
 The kernel processes the buffer as a single contiguous run of ``numel``
 elements; multi-dimensional torch tensors must be ``contiguous()``. This
@@ -51,7 +52,7 @@
 from dataclasses import dataclass
 from typing import Literal, Tuple
 
-from ...core.ir import I32, IRBuilder, KernelDef, PtrType, Value
+from ...core.ir import I32, IRBuilder, KernelDef, PtrType, Type, Value
 from ...helpers.activations import _sigmoid_via_exp2, _tanh_via_exp2
 from ...helpers.distribution import (
     TileDistributionEncoding,
@@ -77,7 +78,11 @@
     "exp2",
 ]
 BinaryOp = Literal["add", "sub", "mul", "max", "min", "swiglu", "geglu"]
-DType = Literal["f16", "bf16"]
+DType = Literal["f16", "bf16", "f64"]
+
+# Ops computed exactly in the I/O type; the exp2-based activations are
+# f32-only, so f64 is limited to this set.
+_F64_OPS = ("copy", "neg", "abs", "relu", "add", "sub", "mul", "max", "min")
 
 
 @dataclass(frozen=True)
@@ -133,11 +138,16 @@ def elems_per_block(self) -> int:
 def is_valid_spec(spec: ElementwiseSpec) -> Tuple[bool, str]:
     if not (spec.is_unary() or spec.is_binary()):
         return False, f"unknown op {spec.op!r}"
-    if spec.dtype not in ("f16", "bf16"):
+    if spec.dtype not in ("f16", "bf16", "f64"):
         return False, f"unsupported dtype {spec.dtype!r}"
+    if spec.dtype == "f64" and spec.op not in _F64_OPS:
+        return False, f"op {spec.op!r} not supported for f64"
     if spec.block_size not in (64, 128, 256, 512, 1024):
         return False, f"block_size {spec.block_size} not in {{64, 128, 256, 512, 1024}}"
-    if spec.vec not in (2, 4, 8):
+    if spec.dtype == "f64":
+        if spec.vec != 2:
+            return False, f"vec {spec.vec} not in {{2}} for f64"
+    elif spec.vec not in (2, 4, 8):
         return False, f"vec {spec.vec} not in {{2, 4, 8}}"
     return True, "ok"
 
@@ -184,7 +194,8 @@ def _apply_unary(b: IRBuilder, x: Value, op: str) -> Value:
         # ``fsub(0.0, x)`` constant load.
         return b.fmax(x, b.fneg(x))
     if op == "relu":
-        return b.fmax(x, b.const_f32(0.0))
+        zero = b.const_f64(0.0) if x.type.name == "f64" else b.const_f32(0.0)
+        return b.fmax(x, zero)
     if op == "exp2":
         return b.exp2(x)
     if op == "tanh":
@@ -235,6 +246,16 @@ def _apply_binary(b: IRBuilder, a: Value, c: Value, op: str) -> Value:
     raise ValueError(f"unsupported binary op {op!r}")
 
 
+def _to_compute(b: IRBuilder, x: Value) -> Value:
+    """Promote an I/O scalar to the compute type (f64 stays native)."""
+    return x if x.type.name == "f64" else b.cast_to_f32(x)
+
+
+def _from_compute(b: IRBuilder, x: Value, io_ty: Type) -> Value:
+    """Demote a compute scalar back to the I/O type (f64 stays native)."""
+    return x if io_ty.name == "f64" else b.cast_f32_to(x, io_ty)
+
+
 # ---------------------------------------------------------------------
 # Codegen
 # ---------------------------------------------------------------------
@@ -351,13 +372,13 @@ def emit_scalar_path() -> None:
             idx = b.add(thread_base, b.const_i32(i))
             in_bounds = b.cmp_lt(idx, N)
             with b.scf_if(in_bounds):
-                a = b.cast_to_f32(a_view.load_scalar(b, [idx]))
+                a = _to_compute(b, a_view.load_scalar(b, [idx]))
                 if spec.is_binary():
-                    bv = b.cast_to_f32(b_view.load_scalar(b, [idx]))
+                    bv = _to_compute(b, b_view.load_scalar(b, [idx]))
                     r = _apply_binary(b, a, bv, spec.op)
                 else:
                     r = _apply_unary(b, a, spec.op)
-                c_view.store_scalar(b, [idx], b.cast_f32_to(r, io_ty))
+                c_view.store_scalar(b, [idx], _from_compute(b, r, io_ty))
 
     with b.scf_if(in_fast):
         emit_vec_path()
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/common/reduce.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/common/reduce.py
index 0b168d728ca4..8facfb4350e2 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/common/reduce.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/common/reduce.py
@@ -10,8 +10,8 @@
     max :  Y[m] = max_n(X[m,n])
     mean:  Y[m] = sum_n(X[m,n]) / N
 
-Compute is in f32 internally; the output dtype matches the input dtype
-(f16 in / f16 out, etc.).
+Compute is in f32 internally (f64 input computes natively in f64); the
+output dtype matches the input dtype (f16 in / f16 out, etc.).
 
 CK Tile parity shape (``BlockReduce2d{Sync, CrossWarpSync}`` in
 :file:`include/ck_tile/ops/reduce/block/block_reduce2d.hpp`):
@@ -40,7 +40,7 @@
 from dataclasses import dataclass
 from typing import Literal, Tuple
 
-from ...core.ir import F32, I32, IRBuilder, KernelDef, PtrType, Value
+from ...core.ir import F32, F64, I32, IRBuilder, KernelDef, PtrType, Value
 from ...helpers.distribution import (
     TileDistributionEncoding,
     block_tile_reduce_sync,
@@ -69,7 +69,7 @@
 )
 
 
-DType = Literal["f16", "bf16"]
+DType = Literal["f16", "bf16", "f64"]
 ReduceOp = Literal["sum", "max", "min", "mean", "prod"]
 
 
@@ -115,12 +115,16 @@ def is_valid_spec(spec: Reduce2DSpec) -> Tuple[bool, str]:
             block_size=spec.block_size,
             vec=spec.vec,
             n_per_block=spec.n_per_block,
+            allowed_dtypes=("f16", "fp16", "bf16", "f64"),
+            allowed_vecs=(2,) if spec.dtype == "f64" else (2, 4, 8),
         )
     )
 
 
 _NEG_INF_F32 = -3.4028234663852886e38
 _POS_INF_F32 = 3.4028234663852886e38
+_NEG_INF_F64 = float("-inf")
+_POS_INF_F64 = float("inf")
 
 
 def _combine_scalar(
@@ -129,7 +133,7 @@ def _combine_scalar(
     a: Value,
     c: Value,
 ) -> Value:
-    """One step of the per-element reduction combiner in f32.
+    """One step of the per-element reduction combiner in the compute type.
 
     Centralised so the per-chunk tree fold, the warp XOR butterfly,
     and the cross-warp LDS tree share the exact same op selection.
@@ -209,20 +213,25 @@ def build_reduce2d(spec: Reduce2DSpec) -> KernelDef:
     x_view = make_naive_tensor_view_packed(X, shape=(1, N), dtype=io_ty)
     x_tile = make_tile_window(x_view, lengths=(1, N), origin=(row, b.const_i32(0)))
 
-    lds = make_lds_view(b, dtype=F32, shape=(BS,), name_hint="lds_red").base
+    # f16/bf16 accumulate in f32; f64 accumulates natively.
+    is_f64 = io_ty == F64
+    compute_ty = F64 if is_f64 else F32
+    const_c = b.const_f64 if is_f64 else b.const_f32
 
-    # Pick the f32 identity element + the combiner for each op.
+    lds = make_lds_view(b, dtype=compute_ty, shape=(BS,), name_hint="lds_red").base
+
+    # Pick the identity element + the combiner for each op.
     if spec.op in ("sum", "mean"):
-        acc = b.const_f32(0.0)
+        acc = const_c(0.0)
         combine: Literal["sum", "max", "min", "prod"] = "sum"
     elif spec.op == "max":
-        acc = b.const_f32(_NEG_INF_F32)
+        acc = const_c(_NEG_INF_F64 if is_f64 else _NEG_INF_F32)
         combine = "max"
     elif spec.op == "min":
-        acc = b.const_f32(_POS_INF_F32)
+        acc = const_c(_POS_INF_F64 if is_f64 else _POS_INF_F32)
         combine = "min"
     elif spec.op == "prod":
-        acc = b.const_f32(1.0)
+        acc = const_c(1.0)
         combine = "prod"
     else:  # validated earlier; defensive
         raise ValueError(f"unsupported reduce op {spec.op!r}")
@@ -270,7 +279,7 @@ def body(_n_off, x_scalars):
     # that isn't a clean multiple of ``wave_size``).
     if spec.block_size % spec.wave_size == 0 and combine in ("sum", "max"):
         red_dist = _make_row_reduce_distribution(spec)
-        reduced = make_static_distributed_tensor(red_dist, dtype=F32)
+        reduced = make_static_distributed_tensor(red_dist, dtype=compute_ty)
         reduced.storage[0] = acc
         block_tile_reduce_sync(
             b,
@@ -297,10 +306,13 @@ def body(_n_off, x_scalars):
         total = block_lds_reduce(b, acc, lds, tid, block_size=BS, combine=combine)
 
     if spec.op == "mean":
-        total = b.fmul(total, b.rcp(b.const_f32(float(N))))
+        total = b.fmul(total, b.rcp(const_c(float(N))))
 
     with b.scf_if(b.cmp_eq(tid, b.const_i32(0))):
-        store_scalar_from_f32(b, Y, row, total, dtype=spec.dtype)
+        if is_f64:
+            b.global_store(Y, row, total)
+        else:
+            store_scalar_from_f32(b, Y, row, total, dtype=spec.dtype)
 
     return b.kernel
 
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_amdgpu_lowering_ops.py b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_amdgpu_lowering_ops.py
index 97df138a2ae9..e940a146ae30 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_amdgpu_lowering_ops.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_amdgpu_lowering_ops.py
@@ -129,9 +129,9 @@ def test_unsupported_elem_type_raises_instead_of_falling_back(self):
         from rocke.core.lower_hip import _vec_prefix
 
         with self.assertRaises(NotImplementedError) as cm:
-            _vec_prefix("f64", "global_load_vN")
+            _vec_prefix("i64", "global_load_vN")
         self.assertIn("global_load_vN", str(cm.exception))
-        self.assertIn("f64", str(cm.exception))
+        self.assertIn("i64", str(cm.exception))
 
 
 class TestHipArchSeamValidation(unittest.TestCase):
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_dtypes.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_dtypes.cpp
index 84e94d2d63e1..5a74081ed011 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_dtypes.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_dtypes.cpp
@@ -13,6 +13,8 @@ int main()
     const char* cases[][2] = {{" HALF ", "fp16"},
                               {"bfloat16", "bf16"},
                               {" Float\t", "fp32"},
+                              {" Double ", "fp64"},
+                              {"F64", "fp64"},
                               {"FP8", "fp8e4m3"},
                               {"BF8", "bf8e5m2"},
                               {"FP6", "fp6e2m3"},
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_dtypes.py b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_dtypes.py
index f5018fbeef28..38db0219d538 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_dtypes.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_dtypes.py
@@ -32,6 +32,7 @@
         ("f16", 16),
         ("bf16", 16),
         ("f32", 32),
+        ("f64", 64),
         ("e8m0", 8),
         ("e4m3", 8),
         ("e5m3", 8),
@@ -72,6 +73,8 @@ def test_unknown_dtype_and_unrepresented_integer_family():
         (" HALF ", "fp16"),
         ("bfloat16", "bf16"),
         (" Float\t", "fp32"),
+        (" Double ", "fp64"),
+        ("F64", "fp64"),
         ("FP8", "fp8e4m3"),
         (" E4M3 ", "fp8e4m3"),
         ("BF8", "bf8e5m2"),
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/numeric.py b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/numeric.py
index 538002d226ba..6486e5f4073a 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/numeric.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/numeric.py
@@ -85,6 +85,9 @@ class Tol:
     "fp32": Tol(rtol=1e-5, atol=1e-6),
     "fp16": Tol(rtol=1e-2, atol=1e-2),
     "bf16": Tol(rtol=1.5e-2, atol=1e-2),
+    # fp64 computes natively; the elementwise inputs are exact in f64, so only
+    # reassociation in reductions contributes error (see _ROW_TOL_FP64).
+    "fp64": Tol(rtol=1e-12, atol=1e-12),
 }
 
 
@@ -299,7 +302,17 @@ def passed(self) -> bool:
 # torch helpers
 # ---------------------------------------------------------------------
 def _torch_dtype(torch, dtype: str):
-    return {"fp16": torch.float16, "bf16": torch.bfloat16, "fp32": torch.float32}[dtype]
+    return {
+        "fp16": torch.float16,
+        "bf16": torch.bfloat16,
+        "fp32": torch.float32,
+        "fp64": torch.float64,
+    }[dtype]
+
+
+def _ref_dtype(torch, td):
+    """Reference compute dtype: float64 for fp64 I/O, float32 otherwise."""
+    return torch.float64 if td == torch.float64 else torch.float32
 
 
 def _make_inputs(torch, m: int, n: int, k: int, dtype: str, seed: int = 0xC0FFEE):
@@ -324,15 +337,17 @@ def _make_inputs(torch, m: int, n: int, k: int, dtype: str, seed: int = 0xC0FFEE
     return A, B
 
 
-def _compare(torch, out, ref_f32, tol: Tol) -> Tuple[float, float, float]:
-    """Return (max_abs_diff, max_rel_diff, worst allclose margin)."""
-    out_f32 = out.to(torch.float32)
-    diff = (out_f32 - ref_f32).abs()
+def _compare(torch, out, ref, tol: Tol) -> Tuple[float, float, float]:
+    """Return (max_abs_diff, max_rel_diff, worst allclose margin).
+
+    ``out`` is compared in ``ref``'s dtype (float32, or float64 for fp64).
+    """
+    diff = (out.to(ref.dtype) - ref).abs()
     max_abs = float(diff.max().item())
-    denom = ref_f32.abs().clamp_min(1e-12)
+    denom = ref.abs().clamp_min(1e-12)
     max_rel = float((diff / denom).max().item())
     # allclose margin: |d| - (atol + rtol*|ref|), worst (max) over elems
-    allowed = tol.atol + tol.rtol * ref_f32.abs()
+    allowed = tol.atol + tol.rtol * ref.abs()
     margin = float((diff - allowed).max().item())
     return max_abs, max_rel, margin
 
@@ -523,7 +538,7 @@ class ElemCfg:
     name: str
     n: int  # numel (kept a multiple of block_size*vec -> fast path, no remainder)
     op: str  # binary: add|sub|mul|max|min   unary: relu|silu|...
-    dtype: str  # "f16" | "bf16"
+    dtype: str  # "f16" | "bf16" | "f64"
     block_size: int = 256
     vec: int = 8
 
@@ -538,25 +553,31 @@ class ElemCfg:
     ElemCfg("add_bf16_128k", 2048 * 64, "add", "bf16"),
     ElemCfg("silu_bf16_128k", 2048 * 64, "silu", "bf16"),
     ElemCfg("tanh_bf16_128k", 2048 * 64, "tanh", "bf16"),
+    # f64: exact ops only, vec=2 (256*2 = 512 elems/block).
+    ElemCfg("add_f64_128k", 2048 * 64, "add", "f64", vec=2),
+    ElemCfg("mul_f64_128k", 2048 * 64, "mul", "f64", vec=2),
+    ElemCfg("min_f64_128k", 2048 * 64, "min", "f64", vec=2),
+    ElemCfg("relu_f64_128k", 2048 * 64, "relu", "f64", vec=2),
 ]
 
 # map spec dtype -> TOL-table key
-_ELEM_TOL_KEY = {"f16": "fp16", "bf16": "bf16"}
+_ELEM_TOL_KEY = {"f16": "fp16", "bf16": "bf16", "f64": "fp64"}
 
 
 def _torch_elementwise_ref(torch, op: str, A, B):
-    """Reference for one elementwise op, computed in fp32."""
-    a = A.to(torch.float32)
+    """Reference for one elementwise op, computed in fp32 (fp64 for fp64 I/O)."""
+    rd = _ref_dtype(torch, A.dtype)
+    a = A.to(rd)
     if op == "add":
-        return a + B.to(torch.float32)
+        return a + B.to(rd)
     if op == "sub":
-        return a - B.to(torch.float32)
+        return a - B.to(rd)
     if op == "mul":
-        return a * B.to(torch.float32)
+        return a * B.to(rd)
     if op == "max":
-        return torch.maximum(a, B.to(torch.float32))
+        return torch.maximum(a, B.to(rd))
     if op == "min":
-        return torch.minimum(a, B.to(torch.float32))
+        return torch.minimum(a, B.to(rd))
     if op == "relu":
         return torch.relu(a)
     if op == "silu":
@@ -661,7 +682,7 @@ def run_elementwise_config(cfg: ElemCfg, arch: str = "gfx950") -> NumericResult:
         ).to(td)
     C = torch.empty((cfg.n,), device="cuda", dtype=td)
 
-    ref_f32 = (
+    ref = (
         reference
         if cfg.op == "tanh"
         else _torch_elementwise_ref(torch, cfg.op, A, B if is_binary else None)
@@ -687,25 +708,25 @@ def run_elementwise_config(cfg: ElemCfg, arch: str = "gfx950") -> NumericResult:
         return res
 
     compare_actual = C
-    compare_reference = ref_f32
+    compare_reference = ref
     if cfg.op == "tanh":
         compare_actual = C.cpu()
         actual_nan = torch.isnan(compare_actual)
-        reference_nan = torch.isnan(ref_f32)
+        reference_nan = torch.isnan(ref)
         if not torch.equal(actual_nan, reference_nan):
             res.status = "DRIFT"
             res.detail = "tanh NaN classification mismatch"
             return res
-        zero_mask = ref_f32 == 0
+        zero_mask = ref == 0
         if not torch.equal(
             torch.signbit(compare_actual[zero_mask]),
-            torch.signbit(ref_f32[zero_mask]),
+            torch.signbit(ref[zero_mask]),
         ):
             res.status = "DRIFT"
             res.detail = "tanh signed-zero mismatch"
             return res
         compare_actual = compare_actual[~reference_nan]
-        compare_reference = ref_f32[~reference_nan]
+        compare_reference = ref[~reference_nan]
 
     max_abs, max_rel, margin = _compare(torch, compare_actual, compare_reference, tol)
     res.max_abs_diff = max_abs
@@ -877,7 +898,7 @@ class RowCfg:
     family: str  # "layernorm2d" | "rmsnorm2d" | "reduce2d"
     m: int
     n: int  # == n_per_block
-    dtype: str  # "f16" | "bf16"
+    dtype: str  # "f16" | "bf16" | "f64" (reduce2d only)
     op: str = ""  # reduce only: sum|max|min|mean
     block_size: int = 256
     vec: int = 8
@@ -897,6 +918,11 @@ class RowCfg:
     RowCfg("red_max_f16_512x4096", "reduce2d", 512, 4096, "f16", op="max"),
     RowCfg("red_mean_f16_512x4096", "reduce2d", 512, 4096, "f16", op="mean"),
     RowCfg("red_sum_bf16_512x4096", "reduce2d", 512, 4096, "bf16", op="sum"),
+    # reduce2d f64 (native f64 accumulate, vec=2); min takes the wave-XOR prologue.
+    RowCfg("red_sum_f64_512x4096", "reduce2d", 512, 4096, "f64", op="sum", vec=2),
+    RowCfg("red_max_f64_512x4096", "reduce2d", 512, 4096, "f64", op="max", vec=2),
+    RowCfg("red_mean_f64_512x4096", "reduce2d", 512, 4096, "f64", op="mean", vec=2),
+    RowCfg("red_min_f64_512x4096", "reduce2d", 512, 4096, "f64", op="min", vec=2),
 ]
 
 # Per-family tolerance overrides. layernorm/rmsnorm carry a divide by a
@@ -908,9 +934,14 @@ class RowCfg:
     # keep an absolute floor proportional to N and a modest rtol.
     "reduce2d": Tol(rtol=5e-2, atol=5e-2),
 }
+# fp64 reduce: only the summation order differs from the reference, so the
+# error is ~N * eps64 * max|partial| (well under 1e-10 for these shapes).
+_ROW_TOL_FP64 = Tol(rtol=1e-10, atol=1e-10)
 
 
 def _row_tol(cfg: RowCfg) -> Tol:
+    if cfg.dtype == "f64":
+        return _ROW_TOL_FP64
     return _ROW_TOL[cfg.family]
 
 
@@ -1040,7 +1071,7 @@ def run_row_config(cfg: RowCfg, arch: str = "gfx950") -> NumericResult:
     torch.manual_seed(0xC0FFEE)
     # small magnitudes so f16/bf16 round-trips and accumulation are tame
     X = (torch.randn((cfg.m, cfg.n), device="cuda", dtype=torch.float32) * 0.5).to(td)
-    x32 = X.to(torch.float32)
+    x_ref = X.to(_ref_dtype(torch, td))
     grid = grid_fn(cfg.m, spec)
     block = (spec.block_size, 1, 1)
     sig = sig_fn(spec)
@@ -1051,10 +1082,10 @@ def run_row_config(cfg: RowCfg, arch: str = "gfx950") -> NumericResult:
         ).to(td)
         Beta = (torch.randn((cfg.n,), device="cuda", dtype=torch.float32) * 0.1).to(td)
         Y = torch.zeros((cfg.m, cfg.n), device="cuda", dtype=td)
-        mean = x32.mean(dim=-1, keepdim=True)
-        var = ((x32 - mean) ** 2).mean(dim=-1, keepdim=True)
+        mean = x_ref.mean(dim=-1, keepdim=True)
+        var = ((x_ref - mean) ** 2).mean(dim=-1, keepdim=True)
         inv_std = 1.0 / torch.sqrt(var + eps)
-        ref_f32 = (x32 - mean) * inv_std * Gamma.to(torch.float32)[None, :] + Beta.to(
+        ref = (x_ref - mean) * inv_std * Gamma.to(torch.float32)[None, :] + Beta.to(
             torch.float32
         )[None, :]
         out = Y
@@ -1072,20 +1103,20 @@ def run_row_config(cfg: RowCfg, arch: str = "gfx950") -> NumericResult:
             torch.randn((cfg.n,), device="cuda", dtype=torch.float32) * 0.1 + 1.0
         ).to(td)
         Y = torch.zeros((cfg.m, cfg.n), device="cuda", dtype=td)
-        rms = torch.sqrt((x32**2).mean(dim=-1, keepdim=True) + eps)
-        ref_f32 = (x32 / rms) * Gamma.to(torch.float32)[None, :]
+        rms = torch.sqrt((x_ref**2).mean(dim=-1, keepdim=True) + eps)
+        ref = (x_ref / rms) * Gamma.to(torch.float32)[None, :]
         out = Y
         values = {"X": X, "Gamma": Gamma, "Y": Y, "M": cfg.m, "N": cfg.n, "eps": eps}
     else:  # reduce2d
         Y = torch.zeros((cfg.m,), device="cuda", dtype=td)
         if cfg.op == "sum":
-            ref_f32 = x32.sum(dim=-1)
+            ref = x_ref.sum(dim=-1)
         elif cfg.op == "max":
-            ref_f32 = x32.max(dim=-1).values
+            ref = x_ref.max(dim=-1).values
         elif cfg.op == "min":
-            ref_f32 = x32.min(dim=-1).values
+            ref = x_ref.min(dim=-1).values
         elif cfg.op == "mean":
-            ref_f32 = x32.mean(dim=-1)
+            ref = x_ref.mean(dim=-1)
         else:  # pragma: no cover
             res.status = "REJECTED"
             res.detail = f"no torch ref for reduce op {cfg.op!r}"
@@ -1104,7 +1135,7 @@ def run_row_config(cfg: RowCfg, arch: str = "gfx950") -> NumericResult:
         res.detail = f"launch raised: {e}"
         return res
 
-    max_abs, max_rel, margin = _compare(torch, out, ref_f32, tol)
+    max_abs, max_rel, margin = _compare(torch, out, ref, tol)
     res.max_abs_diff = max_abs
     res.max_rel_diff = max_rel
     res.margin = margin
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/elementwise_emit.c b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/elementwise_emit.c
index 767b7e756714..a23246730114 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/elementwise_emit.c
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/elementwise_emit.c
@@ -80,6 +80,24 @@ static int make_spec(int idx, rocke_elementwise_spec_t* spec)
         spec->block_size = 256;
         spec->vec = 8;
         break;
+    case 10: /* f64: natively computed, exact ops only (vec=2) */
+        spec->op = "add";
+        spec->dtype = "f64";
+        spec->block_size = 256;
+        spec->vec = 2;
+        break;
+    case 11:
+        spec->op = "relu";
+        spec->dtype = "f64";
+        spec->block_size = 128;
+        spec->vec = 2;
+        break;
+    case 12:
+        spec->op = "max";
+        spec->dtype = "f64";
+        spec->block_size = 512;
+        spec->vec = 2;
+        break;
     default:
         return -1;
     }
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/elementwise_emit.py b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/elementwise_emit.py
index 986f606b63e4..5f352b9945a1 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/elementwise_emit.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/elementwise_emit.py
@@ -44,6 +44,13 @@ def _spec(idx: int) -> ElementwiseSpec:
         return ElementwiseSpec(op="tanh", dtype="f16", block_size=256, vec=8)
     if idx == 9:
         return ElementwiseSpec(op="tanh", dtype="bf16", block_size=256, vec=8)
+    if idx == 10:
+        # f64: natively computed, exact ops only (vec=2).
+        return ElementwiseSpec(op="add", dtype="f64", block_size=256, vec=2)
+    if idx == 11:
+        return ElementwiseSpec(op="relu", dtype="f64", block_size=128, vec=2)
+    if idx == 12:
+        return ElementwiseSpec(op="max", dtype="f64", block_size=512, vec=2)
     raise SystemExit(f"unknown config index {idx}")
 
 
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/reduce_emit.c b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/reduce_emit.c
index 376e9fec47dd..ea0fa34db4e1 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/reduce_emit.c
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/reduce_emit.c
@@ -83,6 +83,46 @@ static int make_spec(int idx, rocke_reduce2d_spec_t* spec)
         spec->dtype = "f16";
         spec->wave_size = 32;
         break;
+    case 8: /* f64: accumulates natively (block_tile_reduce_sync, vec=2) */
+        spec->n_per_block = 4096;
+        spec->op = "sum";
+        spec->block_size = 256;
+        spec->vec = 2;
+        spec->dtype = "f64";
+        spec->wave_size = 64;
+        break;
+    case 9:
+        spec->n_per_block = 2048;
+        spec->op = "max";
+        spec->block_size = 256;
+        spec->vec = 2;
+        spec->dtype = "f64";
+        spec->wave_size = 64;
+        break;
+    case 10:
+        spec->n_per_block = 2048;
+        spec->op = "mean";
+        spec->block_size = 128;
+        spec->vec = 2;
+        spec->dtype = "f64";
+        spec->wave_size = 64;
+        break;
+    case 11: /* f64 min / prod take the wave-XOR prologue path */
+        spec->n_per_block = 4096;
+        spec->op = "min";
+        spec->block_size = 256;
+        spec->vec = 2;
+        spec->dtype = "f64";
+        spec->wave_size = 64;
+        break;
+    case 12:
+        spec->n_per_block = 4096;
+        spec->op = "prod";
+        spec->block_size = 256;
+        spec->vec = 2;
+        spec->dtype = "f64";
+        spec->wave_size = 64;
+        break;
     default:
         return -1;
     }
@@ -103,7 +143,7 @@ int main(int argc, char** argv)
 {
     if(argc < 2)
     {
-        fprintf(stderr, "usage: %s <config_index 0..7> [mode]\n", argv[0]);
+        fprintf(stderr, "usage: %s <config_index 0..12> [mode]\n", argv[0]);
         return 2;
     }
     int idx = atoi(argv[1]);
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/reduce_emit.py b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/reduce_emit.py
index 5b63c9bea59d..77e27711cf3f 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/reduce_emit.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/reduce_emit.py
@@ -80,6 +80,38 @@ def _spec(idx: int) -> Reduce2DSpec:
             ),
             "gfx1201",
         )
+    if idx == 8:
+        # f64: accumulates natively (block_tile_reduce_sync, vec=2).
+        return Reduce2DSpec(
+            n_per_block=4096, op="sum", block_size=256, vec=2, dtype="f64", wave_size=64
+        )
+    if idx == 9:
+        return Reduce2DSpec(
+            n_per_block=2048, op="max", block_size=256, vec=2, dtype="f64", wave_size=64
+        )
+    if idx == 10:
+        return Reduce2DSpec(
+            n_per_block=2048,
+            op="mean",
+            block_size=128,
+            vec=2,
+            dtype="f64",
+            wave_size=64,
+        )
+    if idx == 11:
+        # f64 min / prod take the wave-XOR prologue path.
+        return Reduce2DSpec(
+            n_per_block=4096, op="min", block_size=256, vec=2, dtype="f64", wave_size=64
+        )
+    if idx == 12:
+        return Reduce2DSpec(
+            n_per_block=4096,
+            op="prod",
+            block_size=256,
+            vec=2,
+            dtype="f64",
+            wave_size=64,
+        )
     raise SystemExit(f"unknown config index {idx}")
 
 
@@ -91,7 +123,7 @@ def main() -> int:
     return run_emit(
         _spec,
         _build,
-        usage="usage: reduce_emit.py <config_index 0..7> [mode]\n",
+        usage="usage: reduce_emit.py <config_index 0..12> [mode]\n",
     )
 
 
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py b/dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py
index 3b8d8dd7380b..2d05bb71611e 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py
@@ -1032,6 +1032,22 @@ def test_rejects_unknown_op(self):
         with self.assertRaises(ValueError):
             build_elementwise(ElementwiseSpec(op="bogus"))
 
+    def test_f64_path_builds(self):
+        from rocke.instances import ElementwiseSpec, build_elementwise
+
+        kernel = build_elementwise(ElementwiseSpec(op="add", dtype="f64", vec=2))
+        ll = lower_kernel_to_llvm(kernel)
+        self.assertIn("fadd double", ll)
+
+    def test_f64_rejects_transcendental_ops_and_wide_vec(self):
+        from rocke.instances.common.elementwise import ElementwiseSpec, is_valid_spec
+
+        for op in ("exp2", "tanh", "silu", "gelu_tanh"):
+            ok, _ = is_valid_spec(ElementwiseSpec(op=op, dtype="f64", vec=2))
+            self.assertFalse(ok, op)
+        ok, _ = is_valid_spec(ElementwiseSpec(op="add", dtype="f64", vec=4))
+        self.assertFalse(ok)
+
 
 class TestLayerNormInstance(unittest.TestCase):
     def test_builds_with_save_mean(self):
@@ -1070,6 +1086,31 @@ def test_sum_max_mean_builds(self):
             ll = lower_kernel_to_llvm(kernel)
             self.assertIn("define amdgpu_kernel void", ll)
 
+    def test_f64_path_builds(self):
+        from rocke.instances import Reduce2DSpec, build_reduce2d
+
+        for op in ("sum", "max", "mean"):
+            spec = Reduce2DSpec(n_per_block=4096, op=op, dtype="f64", vec=2)
+            ll = lower_kernel_to_llvm(build_reduce2d(spec))
+            self.assertIn("double", ll)
+
+    def test_f64_max_min_identities_are_infinite(self):
+        from rocke.instances import Reduce2DSpec, build_reduce2d
+
+        for op, inf_bits in (
+            ("max", "0xFFF0000000000000"),
+            ("min", "0x7FF0000000000000"),
+        ):
+            spec = Reduce2DSpec(n_per_block=4096, op=op, dtype="f64", vec=2)
+            ll = lower_kernel_to_llvm(build_reduce2d(spec))
+            self.assertIn(f"double {inf_bits}", ll)
+
+    def test_f64_rejects_wide_vec(self):
+        from rocke.instances.common.reduce import Reduce2DSpec, is_valid_spec
+
+        ok, _ = is_valid_spec(Reduce2DSpec(n_per_block=4096, dtype="f64", vec=4))
+        self.assertFalse(ok)
+
 
 class TestTransposeInstance(unittest.TestCase):
     def test_builds(self):

From 749420b65c4cbaf2d1f0f93f0a34bbe2dbbc7e7d Mon Sep 17 00:00:00 2001
From: Minseob Shin <minseob.shin11@gmail.com>
Date: Wed, 7 Oct 2026 20:00:08 +0000
Subject: [PATCH 2/7] change f64 ordering and test input

---
 .../rocke/platform/cpp/core/ir/ir_mem.cpp     | 28 +++++++-------
 .../rocke/platform/cpp/core/ir/ir_tile.cpp    |  8 ++--
 .../platform/cpp/core/lower_hip/core.cpp      |  4 +-
 .../platform/cpp/core/lower_llvm/core.cpp     | 10 ++---
 .../platform/cpp/core/lower_llvm/data.cpp     |  4 +-
 .../cpp/include/rocke/instance_elementwise.h  |  2 +-
 .../rocke/platform/cpp/include/rocke/ir.h     |  2 +-
 .../cpp/include/rocke/lower_llvm_internal.h   |  6 +--
 .../cpp/instances/common/elementwise.cpp      |  4 +-
 .../platform/dsl_docs/ir_lowering/ir_model.md |  2 +-
 .../rocke/platform/python/rocke/core/ir.py    | 38 +++++++++----------
 .../platform/python/rocke/core/lower_llvm.py  | 32 ++++++++--------
 .../python/rocke/instances/SUPPORT_MATRIX.md  |  2 +-
 .../rocke/instances/common/elementwise.py     |  6 +--
 .../tests/instances/differential/numeric.py   | 31 +++++++++------
 .../tests/instances/parity/elementwise_emit.c |  2 +-
 .../instances/parity/elementwise_emit.py      |  2 +-
 17 files changed, 95 insertions(+), 88 deletions(-)

diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_mem.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_mem.cpp
index 268c913a3234..5844651865c1 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_mem.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_mem.cpp
@@ -290,18 +290,18 @@ rocke_value_t* rocke_b_global_load_vN_ex(rocke_ir_builder_t* b,
             return (rocke_value_t*)rocke_i_set_err(
                 b, ROCKE_ERR_VALUE, "unsupported vector width for %s global_load_vN: %d", en, n);
     }
-    else if(rocke_i_type_is(dtype, "fp8e4m3") || rocke_i_type_is(dtype, "bf8e5m2")
-            || rocke_i_type_is(dtype, "i8"))
+    else if(rocke_i_type_is(dtype, "f64"))
     {
-        elem_bytes = 1;
-        if(n != 2 && n != 4 && n != 8 && n != 12 && n != 16)
+        elem_bytes = 8;
+        if(n != 2)
             return (rocke_value_t*)rocke_i_set_err(
                 b, ROCKE_ERR_VALUE, "unsupported vector width for %s global_load_vN: %d", en, n);
     }
-    else if(rocke_i_type_is(dtype, "f64"))
+    else if(rocke_i_type_is(dtype, "fp8e4m3") || rocke_i_type_is(dtype, "bf8e5m2")
+            || rocke_i_type_is(dtype, "i8"))
     {
-        elem_bytes = 8;
-        if(n != 2)
+        elem_bytes = 1;
+        if(n != 2 && n != 4 && n != 8 && n != 12 && n != 16)
             return (rocke_value_t*)rocke_i_set_err(
                 b, ROCKE_ERR_VALUE, "unsupported vector width for %s global_load_vN: %d", en, n);
     }
@@ -310,7 +310,7 @@ rocke_value_t* rocke_b_global_load_vN_ex(rocke_ir_builder_t* b,
         return (rocke_value_t*)rocke_i_set_err(
             b,
             ROCKE_ERR_VALUE,
-            "global_load_vN supports f16/bf16/i16/f32/i32/tf32/fp8e4m3/bf8e5m2/i8/f64, got %s",
+            "global_load_vN supports f16/bf16/i16/f32/i32/tf32/f64/fp8e4m3/bf8e5m2/i8, got %s",
             en);
     }
     streaming = mem_opts_streaming(b, opts, "global_load_vN");
@@ -426,11 +426,6 @@ void rocke_b_global_store_vN_ex(rocke_ir_builder_t* b,
             return;
         }
     }
-    else if(rocke_i_type_is(et, "i8") || rocke_i_type_is(et, "fp8e4m3")
-            || rocke_i_type_is(et, "bf8e5m2"))
-    {
-        elem_bytes = 1;
-    }
     else if(rocke_i_type_is(et, "f64"))
     {
         elem_bytes = 8;
@@ -441,12 +436,17 @@ void rocke_b_global_store_vN_ex(rocke_ir_builder_t* b,
             return;
         }
     }
+    else if(rocke_i_type_is(et, "i8") || rocke_i_type_is(et, "fp8e4m3")
+            || rocke_i_type_is(et, "bf8e5m2"))
+    {
+        elem_bytes = 1;
+    }
     else
     {
         (void)rocke_i_set_err(
             b,
             ROCKE_ERR_VALUE,
-            "global_store_vN supports f16/bf16/i16/f32/i32/tf32/i8/fp8e4m3/bf8e5m2/f64, got %s",
+            "global_store_vN supports f16/bf16/i16/f32/i32/tf32/f64/i8/fp8e4m3/bf8e5m2, got %s",
             en);
         return;
     }
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_tile.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_tile.cpp
index 4e7720febfda..540a27a2f754 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_tile.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_tile.cpp
@@ -380,14 +380,14 @@ rocke_value_t* rocke_b_smem_load_vN(rocke_ir_builder_t* b,
     }
     dn = dtype->name;
     if(!(strcmp(dn, "f16") == 0 || strcmp(dn, "bf16") == 0 || strcmp(dn, "f32") == 0
-         || strcmp(dn, "tf32") == 0 || strcmp(dn, "i32") == 0 || strcmp(dn, "fp8e4m3") == 0
-         || strcmp(dn, "bf8e5m2") == 0 || strcmp(dn, "i8") == 0 || strcmp(dn, "f64") == 0))
+         || strcmp(dn, "tf32") == 0 || strcmp(dn, "i32") == 0 || strcmp(dn, "f64") == 0
+         || strcmp(dn, "fp8e4m3") == 0 || strcmp(dn, "bf8e5m2") == 0 || strcmp(dn, "i8") == 0))
     {
         return (rocke_value_t*)rocke_i_set_err(
             b,
             ROCKE_ERR_VALUE,
-            "smem_load_vN supports f16 / bf16 / f32 / i32 / tf32 / fp8e4m3 / "
-            "bf8e5m2 / i8 / f64, got %s",
+            "smem_load_vN supports f16 / bf16 / f32 / i32 / tf32 / f64 / "
+            "fp8e4m3 / bf8e5m2 / i8, got %s",
             dn);
     }
     {
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/core.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/core.cpp
index 1ed97aabb759..2e864d78144e 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/core.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/core.cpp
@@ -548,8 +548,8 @@ const char* rocke_h_f64_literal(rocke_h_lowerer_t* lw, double val)
         out = rocke_arena_strdup(&lw->b->arena, "-0.0");
         return out ? out : "";
     }
-    /* %.17g round-trips every finite double. The text may differ from CPython
-     * repr (same port hazard as rocke_h_f32_literal). */
+    /* NOTE(port): Python emits repr(float); %.17g is a close approximation but
+     * NOT guaranteed byte-identical to CPython repr. Known port hazard. */
     out = rocke_arena_printf(&lw->b->arena, "%.17g", val);
     return out ? out : "";
 }
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp
index 9e025ce0b5c4..33baa68e5f86 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp
@@ -720,7 +720,7 @@ const char* rocke_ll_smem_storage_type(rocke_lower_t* L, const rocke_type_t* sme
 }
 
 /* ====================================================================== */
-/* FP hex constants (Python _fp32_hex / _fp16_hex / _fp64_hex)            */
+/* FP hex constants (Python _fp32_hex / _fp64_hex / _fp16_hex)            */
 /* ====================================================================== */
 
 const char* rocke_ll_fp32_hex(rocke_lower_t* L, double x)
@@ -898,17 +898,17 @@ const char* rocke_ll_operand(rocke_lower_t* L, const rocke_value_t* v)
             rocke_attr_get_float(&op->attrs, "value", &fv);
             return rocke_ll_fp32_hex(L, fv);
         }
-        if(strcmp(ity, "f16") == 0)
+        if(strcmp(ity, "f64") == 0)
         {
             double fv = 0.0;
             rocke_attr_get_float(&op->attrs, "value", &fv);
-            return rocke_ll_fp16_hex(L, fv);
+            return rocke_ll_fp64_hex(L, fv);
         }
-        if(strcmp(ity, "f64") == 0)
+        if(strcmp(ity, "f16") == 0)
         {
             double fv = 0.0;
             rocke_attr_get_float(&op->attrs, "value", &fv);
-            return rocke_ll_fp64_hex(L, fv);
+            return rocke_ll_fp16_hex(L, fv);
         }
         int64_t iv = 0;
         if(!rocke_attr_get_int(&op->attrs, "value", &iv))
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/data.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/data.cpp
index 8f87200c6956..235b1ad66c14 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/data.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/data.cpp
@@ -151,13 +151,13 @@ const rocke_ll_decl_t ROCKE_LL_INTRINSIC_DECLS[] = {
     {"rsqrt.f32", "declare float @llvm.amdgcn.rsq.f32(float)"},
     {"rcp.f32", "declare float @llvm.amdgcn.rcp.f32(float)"},
     {"maxnum.f32", "declare float @llvm.maxnum.f32(float, float)"},
+    {"maxnum.f64", "declare double @llvm.maxnum.f64(double, double)"},
     {"maxnum.f16", "declare half @llvm.maxnum.f16(half, half)"},
     {"maxnum.bf16", "declare bfloat @llvm.maxnum.bf16(bfloat, bfloat)"},
-    {"maxnum.f64", "declare double @llvm.maxnum.f64(double, double)"},
     {"minnum.f32", "declare float @llvm.minnum.f32(float, float)"},
+    {"minnum.f64", "declare double @llvm.minnum.f64(double, double)"},
     {"minnum.f16", "declare half @llvm.minnum.f16(half, half)"},
     {"minnum.bf16", "declare bfloat @llvm.minnum.bf16(bfloat, bfloat)"},
-    {"minnum.f64", "declare double @llvm.minnum.f64(double, double)"},
     {"fabs.f32", "declare float @llvm.fabs.f32(float)"},
     {"fabs.f16", "declare half @llvm.fabs.f16(half)"},
     {"fabs.bf16", "declare bfloat @llvm.fabs.bf16(bfloat)"},
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_elementwise.h b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_elementwise.h
index 7fc8a405be28..b147269d847c 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_elementwise.h
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_elementwise.h
@@ -8,7 +8,7 @@
  * Emits a single AMDGPU kernel that walks one contiguous N-element tensor with
  * vectorised global loads/stores and applies a fused unary or binary operation
  * per element. Compute is f32 internally for f16 or bf16 I/O; f64 I/O computes
- * natively (exact ops only, vec=2).
+ * in f64 (native f64 ops only, vec=2).
  *
  *   Python (elementwise.py)               C99 (this header)
  *   -----------------------------------   --------------------------------------
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/ir.h b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/ir.h
index 9589acfb042e..d795d14b68fa 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/ir.h
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/ir.h
@@ -103,7 +103,6 @@ typedef enum rocke_scalar_kind
     ROCKE_SCALAR_BF16,
     ROCKE_SCALAR_F16,
     ROCKE_SCALAR_F32,
-    ROCKE_SCALAR_F64,
     ROCKE_SCALAR_FP8E4M3,
     ROCKE_SCALAR_BF8E5M2,
     ROCKE_SCALAR_FP4E2M1,
@@ -112,6 +111,7 @@ typedef enum rocke_scalar_kind
     ROCKE_SCALAR_E8M0,
     ROCKE_SCALAR_E5M3,
     ROCKE_SCALAR_TF32,
+    ROCKE_SCALAR_F64,
     ROCKE_SCALAR__COUNT
 } rocke_scalar_kind_t;
 
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_llvm_internal.h b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_llvm_internal.h
index 75280cfe9f17..149282e3579d 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_llvm_internal.h
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_llvm_internal.h
@@ -565,11 +565,11 @@ bool rocke_ll_is_constant(const rocke_value_t* v);
  * + returns 0 if not a constant. */
 int64_t rocke_ll_eval_constant(rocke_lower_t* L, const rocke_value_t* v);
 
-/* fp32 / fp16 LLVM hex constant spellings (Python _fp32_hex / _fp16_hex).
- * Arena-owned. */
+/* fp32 / fp64 / fp16 LLVM hex constant spellings (Python _fp32_hex / _fp64_hex /
+ * _fp16_hex). Arena-owned. */
 const char* rocke_ll_fp32_hex(rocke_lower_t* L, double x);
-const char* rocke_ll_fp16_hex(rocke_lower_t* L, double x);
 const char* rocke_ll_fp64_hex(rocke_lower_t* L, double x);
+const char* rocke_ll_fp16_hex(rocke_lower_t* L, double x);
 
 /* Escape a string for an LLVM asm/string literal (Python
  * _escape_llvm_asm_string): printable ASCII verbatim, else \XX hex. */
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/elementwise.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/elementwise.cpp
index c11fd5b1389c..fe1dbec08f3e 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/elementwise.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/elementwise.cpp
@@ -139,8 +139,8 @@ static bool rocke_ew_reason(char* reason, size_t cap, const char* msg)
     return false;
 }
 
-/* _F64_OPS: ops computed exactly in the I/O type; the exp2-based activations
- * are f32-only, so f64 is limited to this set. */
+/* _F64_OPS: ops with a native f64 lowering; the exp2-based activations are
+ * f32-only, so f64 is limited to this set. */
 static bool rocke_ew_is_f64_op(const char* op)
 {
     return rocke_ew_streq(op, "copy") || rocke_ew_streq(op, "neg") || rocke_ew_streq(op, "abs")
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/ir_model.md b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/ir_model.md
index f6ae39dc9af7..3b587f6c6033 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/ir_model.md
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/ir_model.md
@@ -100,7 +100,7 @@ zero_vec(elem, n)          # f32 / f16 / bf16
 ```text
 add, sub, mul, div, mod                    # integer arithmetic
 fadd, fsub, fmul, fdiv, fneg               # f32 arithmetic
-fmax, fmin                                 # llvm.maxnum / llvm.minnum (f32/f16/bf16/f64)
+fmax, fmin                                 # llvm.maxnum / llvm.minnum (f32/f64/f16/bf16)
 fcmp(pred, a, b)                           # pred in {olt,ole,ogt,oge,oeq,one,ord,uno}
 cmp_lt, cmp_le, cmp_gt, cmp_ge, cmp_eq, cmp_ne   # integer compare (-> i1)
 land, lor, lnot                            # bitwise + i1 logic
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/ir.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/ir.py
index e5d4e330286e..25c063475edb 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/ir.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/ir.py
@@ -1787,10 +1787,10 @@ def global_load_vN(
         """Vectorised global load of N consecutive values.
 
         Supports f16/bf16/i16 (N in {2, 4, 6, 8, 16}), f32/i32/tf32
-        (N in {2, 3, 4, 8}), fp8e4m3/bf8e5m2/i8 (N in {2, 4, 8, 12, 16}), and
-        f64 (N = 2). Loads exactly N elements. Instruction selection depends on
-        target and alignment; 96-bit payloads do not require a 96-bit scalar
-        type.
+        (N in {2, 3, 4, 8}), f64 (N = 2), and fp8e4m3/bf8e5m2/i8
+        (N in {2, 4, 8, 12, 16}). Loads exactly N elements. Instruction
+        selection depends on target and alignment; 96-bit payloads do not
+        require a 96-bit scalar type.
 
         Default alignment is the payload size for power-of-two loads, and
         element alignment for 12-byte loads. An explicit alignment is a caller
@@ -1819,21 +1819,21 @@ def global_load_vN(
                 raise ValueError(
                     f"unsupported vector width for {dtype.name} global_load_vN: {n}"
                 )
-        elif dtype.name in ("fp8e4m3", "bf8e5m2", "i8"):
-            elem_bytes = 1
-            if n not in (2, 4, 8, 12, 16):
-                raise ValueError(
-                    f"unsupported vector width for {dtype.name} global_load_vN: {n}"
-                )
         elif dtype.name == "f64":
             elem_bytes = 8
             if n != 2:
                 raise ValueError(
                     f"unsupported vector width for {dtype.name} global_load_vN: {n}"
                 )
+        elif dtype.name in ("fp8e4m3", "bf8e5m2", "i8"):
+            elem_bytes = 1
+            if n not in (2, 4, 8, 12, 16):
+                raise ValueError(
+                    f"unsupported vector width for {dtype.name} global_load_vN: {n}"
+                )
         else:
             raise ValueError(
-                "global_load_vN supports f16/bf16/i16/f32/i32/tf32/fp8e4m3/bf8e5m2/i8/f64, "
+                "global_load_vN supports f16/bf16/i16/f32/i32/tf32/f64/fp8e4m3/bf8e5m2/i8, "
                 f"got {dtype.name}"
             )
         attrs = {
@@ -2067,14 +2067,14 @@ def smem_load_vN(self, smem: Value, *indices, dtype: Type, n: int = 0) -> Value:
             "f32",
             "i32",
             "tf32",
+            "f64",
             "fp8e4m3",
             "bf8e5m2",
             "i8",
-            "f64",
         ):
             raise ValueError(
-                "smem_load_vN supports f16 / bf16 / f32 / i32 / tf32 / fp8e4m3 / "
-                f"bf8e5m2 / i8 / f64, got {dtype.name}"
+                "smem_load_vN supports f16 / bf16 / f32 / i32 / tf32 / f64 / "
+                f"fp8e4m3 / bf8e5m2 / i8, got {dtype.name}"
             )
         if dtype.name in ("fp8e4m3", "bf8e5m2", "i8"):
             allowed_n = (1, 2, 4, 8, 12, 16)
@@ -4471,8 +4471,8 @@ def global_store_vN(
 
         Supports the full element-type catalog the LLVM lowering already
         emits: ``f16`` / ``bf16`` / ``i16`` (2-byte), ``f32`` / ``i32`` / ``tf32``
-        (4-byte), ``i8`` / ``fp8e4m3`` / ``bf8e5m2`` (1-byte), ``f64``
-        (8-byte, N in {1, 2}). Lowers to
+        (4-byte), ``f64`` (8-byte, N in {1, 2}), ``i8`` / ``fp8e4m3`` /
+        ``bf8e5m2`` (1-byte). Lowers to
         a single ``store <N x elem>`` with the supplied address alignment.
         Payload width and address alignment are independent; target and
         alignment determine whether the transfer uses one machine instruction.
@@ -4498,15 +4498,15 @@ def global_store_vN(
             elem_bytes = 4
             if n == 16:
                 raise ValueError(f"global_store_vN n=16 not supported for {elem_name}")
-        elif elem_name in ("i8", "fp8e4m3", "bf8e5m2"):
-            elem_bytes = 1
         elif elem_name == "f64":
             elem_bytes = 8
             if n > 2:
                 raise ValueError(f"global_store_vN n={n} not supported for {elem_name}")
+        elif elem_name in ("i8", "fp8e4m3", "bf8e5m2"):
+            elem_bytes = 1
         else:
             raise ValueError(
-                "global_store_vN supports f16/bf16/i16/f32/i32/tf32/i8/fp8e4m3/bf8e5m2/f64, "
+                "global_store_vN supports f16/bf16/i16/f32/i32/tf32/f64/i8/fp8e4m3/bf8e5m2, "
                 f"got {elem_name}"
             )
         attrs = {
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py
index 7d1660ea3eaf..2cc4f2d7a6ab 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py
@@ -417,13 +417,13 @@ def _resolve_llvm_flavor() -> str:
     "rsqrt.f32": "declare float @llvm.amdgcn.rsq.f32(float)",
     "rcp.f32": "declare float @llvm.amdgcn.rcp.f32(float)",
     "maxnum.f32": "declare float @llvm.maxnum.f32(float, float)",
+    "maxnum.f64": "declare double @llvm.maxnum.f64(double, double)",
     "maxnum.f16": "declare half @llvm.maxnum.f16(half, half)",
     "maxnum.bf16": "declare bfloat @llvm.maxnum.bf16(bfloat, bfloat)",
-    "maxnum.f64": "declare double @llvm.maxnum.f64(double, double)",
     "minnum.f32": "declare float @llvm.minnum.f32(float, float)",
+    "minnum.f64": "declare double @llvm.minnum.f64(double, double)",
     "minnum.f16": "declare half @llvm.minnum.f16(half, half)",
     "minnum.bf16": "declare bfloat @llvm.minnum.bf16(bfloat, bfloat)",
-    "minnum.f64": "declare double @llvm.minnum.f64(double, double)",
     "fabs.f32": "declare float @llvm.fabs.f32(float)",
     "fabs.f16": "declare half @llvm.fabs.f16(half)",
     "fabs.bf16": "declare bfloat @llvm.fabs.bf16(bfloat)",
@@ -1567,10 +1567,10 @@ def _operand(self, v: Value) -> str:
             val = op.attrs["value"]
             if ity == "f32":
                 return _fp32_hex(val)
-            if ity == "f16":
-                return _fp16_hex(val)
             if ity == "f64":
                 return _fp64_hex(val)
+            if ity == "f16":
+                return _fp16_hex(val)
             return str(int(val))
         return v.name
 
@@ -2249,15 +2249,15 @@ def _op_arith_fmax(self, op: Op) -> None:
         # silently mis-typed half operands as float and broke comgr.
         llvm_ty = {
             "f32": "float",
+            "f64": "double",
             "f16": "half",
             "bf16": "bfloat",
-            "f64": "double",
         }.get(ty_name)
         intrin_key = {
             "f32": "maxnum.f32",
+            "f64": "maxnum.f64",
             "f16": "maxnum.f16",
             "bf16": "maxnum.bf16",
-            "f64": "maxnum.f64",
         }.get(ty_name)
         if llvm_ty is None or intrin_key is None:
             raise NotImplementedError(f"fmax: unsupported FP type {ty_name!r}")
@@ -2272,15 +2272,15 @@ def _op_arith_fmin(self, op: Op) -> None:
         ty_name = a.type.name
         llvm_ty = {
             "f32": "float",
+            "f64": "double",
             "f16": "half",
             "bf16": "bfloat",
-            "f64": "double",
         }.get(ty_name)
         intrin_key = {
             "f32": "minnum.f32",
+            "f64": "minnum.f64",
             "f16": "minnum.f16",
             "bf16": "minnum.bf16",
-            "f64": "minnum.f64",
         }.get(ty_name)
         if llvm_ty is None or intrin_key is None:
             raise NotImplementedError(f"fmin: unsupported FP type {ty_name!r}")
@@ -6424,6 +6424,14 @@ def _fp32_hex(x: float) -> str:
     return f"0x{bits:016X}"
 
 
+def _fp64_hex(x: float) -> str:
+    # LLVM textual IR spells `double` hex constants as the raw IEEE-754 bits.
+    import struct
+
+    bits = struct.unpack("<Q", struct.pack("<d", float(x)))[0]
+    return f"0x{bits:016X}"
+
+
 def _fp16_hex(x: float) -> str:
     # LLVM IR accepts fp16 constants as `half 0xH<4 hex digits>`. We use
     # a numerically-correct rounding via the struct module.
@@ -6433,14 +6441,6 @@ def _fp16_hex(x: float) -> str:
     return f"0xH{bits:04X}"
 
 
-def _fp64_hex(x: float) -> str:
-    # LLVM textual IR spells `double` hex constants as the raw IEEE-754 bits.
-    import struct
-
-    bits = struct.unpack("<Q", struct.pack("<d", float(x)))[0]
-    return f"0x{bits:016X}"
-
-
 def _lower_kernel_to_llvm_python(
     kernel: KernelDef,
     *,
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/SUPPORT_MATRIX.md b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/SUPPORT_MATRIX.md
index 434ef8a313ad..a3a046964863 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/SUPPORT_MATRIX.md
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/SUPPORT_MATRIX.md
@@ -41,7 +41,7 @@ atom (CDNA: `pipeline=mem`/`compv*`; gfx1151: `pipeline=mem`, `epilogue=default`
 These emit generic AMDGPU IR; arch only sets the comgr target triple.
 
 `elementwise` and `reduce2d` also accept `dtype="f64"` (native f64
-compute, `vec=2`). f64 elementwise is limited to the exact ops `copy`,
+compute, `vec=2`). f64 elementwise is limited to the native f64 ops `copy`,
 `neg`, `abs`, `relu`, `add`, `sub`, `mul`, `max` and `min`.
 
 ---
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/common/elementwise.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/common/elementwise.py
index 17350724c7fb..5bdbee520d25 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/common/elementwise.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/common/elementwise.py
@@ -14,7 +14,7 @@
   ``tanh``
 * Binary ops: ``add``, ``sub``, ``mul``, ``max``, ``min``
 * Dtypes: ``f16`` and ``bf16`` for I/O (compute is f32 internally), and
-  ``f64`` (computed natively; exact ops only, ``vec=2``)
+  ``f64`` (native f64 ops only, ``vec=2``)
 
 The kernel processes the buffer as a single contiguous run of ``numel``
 elements; multi-dimensional torch tensors must be ``contiguous()``. This
@@ -80,8 +80,8 @@
 BinaryOp = Literal["add", "sub", "mul", "max", "min", "swiglu", "geglu"]
 DType = Literal["f16", "bf16", "f64"]
 
-# Ops computed exactly in the I/O type; the exp2-based activations are
-# f32-only, so f64 is limited to this set.
+# Ops with a native f64 lowering; the exp2-based activations are f32-only,
+# so f64 is limited to this set.
 _F64_OPS = ("copy", "neg", "abs", "relu", "add", "sub", "mul", "max", "min")
 
 
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/numeric.py b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/numeric.py
index 6486e5f4073a..724a1df05ecc 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/numeric.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/numeric.py
@@ -85,8 +85,8 @@ class Tol:
     "fp32": Tol(rtol=1e-5, atol=1e-6),
     "fp16": Tol(rtol=1e-2, atol=1e-2),
     "bf16": Tol(rtol=1.5e-2, atol=1e-2),
-    # fp64 computes natively; the elementwise inputs are exact in f64, so only
-    # reassociation in reductions contributes error (see _ROW_TOL_FP64).
+    # fp64 computes natively; each elementwise op rounds like the reference, so
+    # only reassociation in reductions contributes error (see _ROW_TOL_FP64).
     "fp64": Tol(rtol=1e-12, atol=1e-12),
 }
 
@@ -553,7 +553,7 @@ class ElemCfg:
     ElemCfg("add_bf16_128k", 2048 * 64, "add", "bf16"),
     ElemCfg("silu_bf16_128k", 2048 * 64, "silu", "bf16"),
     ElemCfg("tanh_bf16_128k", 2048 * 64, "tanh", "bf16"),
-    # f64: exact ops only, vec=2 (256*2 = 512 elems/block).
+    # f64: native f64 ops only, vec=2 (256*2 = 512 elems/block).
     ElemCfg("add_f64_128k", 2048 * 64, "add", "f64", vec=2),
     ElemCfg("mul_f64_128k", 2048 * 64, "mul", "f64", vec=2),
     ElemCfg("min_f64_128k", 2048 * 64, "min", "f64", vec=2),
@@ -644,12 +644,15 @@ def run_elementwise_config(cfg: ElemCfg, arch: str = "gfx950") -> NumericResult:
     td = _torch_dtype(torch, tol_key)
     torch.manual_seed(0xC0FFEE)
     input_device = "cpu" if cfg.op == "tanh" else "cuda"
-    A = (
-        torch.randint(-4, 5, (cfg.n,), device=input_device, dtype=torch.int32).to(
-            torch.float32
-        )
-        * 0.5
-    ).to(td)
+    if td == torch.float64:
+        A = torch.randn((cfg.n,), device=input_device, dtype=td)
+    else:
+        A = (
+            torch.randint(-4, 5, (cfg.n,), device=input_device, dtype=torch.int32).to(
+                torch.float32
+            )
+            * 0.5
+        ).to(td)
     if cfg.op == "tanh":
         edge_values = torch.tensor(
             [
@@ -673,7 +676,9 @@ def run_elementwise_config(cfg: ElemCfg, arch: str = "gfx950") -> NumericResult:
         A = A.to(device="cuda")
     is_binary = spec.is_binary()
     B = None
-    if is_binary:
+    if is_binary and td == torch.float64:
+        B = torch.randn((cfg.n,), device="cuda", dtype=td)
+    elif is_binary:
         B = (
             torch.randint(-4, 5, (cfg.n,), device="cuda", dtype=torch.int32).to(
                 torch.float32
@@ -1067,10 +1072,12 @@ def run_row_config(cfg: RowCfg, arch: str = "gfx950") -> NumericResult:
     res.extra["kernel_name"] = art.kernel_name
     res.extra["hsaco_bytes"] = art.hsaco_bytes
 
-    # --- deterministic inputs + torch reference (all in f32) ---
+    # --- deterministic inputs + torch reference (f32, or f64 for fp64) ---
     torch.manual_seed(0xC0FFEE)
     # small magnitudes so f16/bf16 round-trips and accumulation are tame
-    X = (torch.randn((cfg.m, cfg.n), device="cuda", dtype=torch.float32) * 0.5).to(td)
+    X = (
+        torch.randn((cfg.m, cfg.n), device="cuda", dtype=_ref_dtype(torch, td)) * 0.5
+    ).to(td)
     x_ref = X.to(_ref_dtype(torch, td))
     grid = grid_fn(cfg.m, spec)
     block = (spec.block_size, 1, 1)
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/elementwise_emit.c b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/elementwise_emit.c
index a23246730114..d53535951ca0 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/elementwise_emit.c
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/elementwise_emit.c
@@ -80,7 +80,7 @@ static int make_spec(int idx, rocke_elementwise_spec_t* spec)
         spec->block_size = 256;
         spec->vec = 8;
         break;
-    case 10: /* f64: natively computed, exact ops only (vec=2) */
+    case 10: /* f64: native f64 ops only (vec=2) */
         spec->op = "add";
         spec->dtype = "f64";
         spec->block_size = 256;
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/elementwise_emit.py b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/elementwise_emit.py
index 5f352b9945a1..5a1bf355c0bf 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/elementwise_emit.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/elementwise_emit.py
@@ -45,7 +45,7 @@ def _spec(idx: int) -> ElementwiseSpec:
     if idx == 9:
         return ElementwiseSpec(op="tanh", dtype="bf16", block_size=256, vec=8)
     if idx == 10:
-        # f64: natively computed, exact ops only (vec=2).
+        # f64: native f64 ops only (vec=2).
         return ElementwiseSpec(op="add", dtype="f64", block_size=256, vec=2)
     if idx == 11:
         return ElementwiseSpec(op="relu", dtype="f64", block_size=128, vec=2)

From 5c3e9646e944c45bdec98014d03b10dd9f5faf3c Mon Sep 17 00:00:00 2001
From: Minseob Shin <minseob.shin11@gmail.com>
Date: Wed, 7 Oct 2026 20:27:40 +0000
Subject: [PATCH 3/7] fix ordering

---
 .../platform/cpp/core/lower_llvm/core.cpp     | 26 +++++++--------
 .../platform/cpp/core/lower_llvm/data.cpp     |  4 +--
 .../cpp/include/rocke/lower_llvm_internal.h   |  4 +--
 .../platform/dsl_docs/ir_lowering/ir_model.md |  2 +-
 .../platform/python/rocke/core/lower_llvm.py  | 32 +++++++++----------
 .../python/rocke/helpers/distribution.py      |  4 +--
 6 files changed, 36 insertions(+), 36 deletions(-)

diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp
index 33baa68e5f86..ae6ddb2aab32 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp
@@ -720,9 +720,17 @@ const char* rocke_ll_smem_storage_type(rocke_lower_t* L, const rocke_type_t* sme
 }
 
 /* ====================================================================== */
-/* FP hex constants (Python _fp32_hex / _fp64_hex / _fp16_hex)            */
+/* FP hex constants (Python _fp64_hex / _fp32_hex / _fp16_hex)            */
 /* ====================================================================== */
 
+const char* rocke_ll_fp64_hex(rocke_lower_t* L, double x)
+{
+    /* LLVM textual IR spells `double` hex constants as the raw IEEE-754 bits. */
+    uint64_t bits;
+    memcpy(&bits, &x, sizeof bits);
+    return rocke_arena_printf(&L->arena, "0x%016llX", (unsigned long long)bits);
+}
+
 const char* rocke_ll_fp32_hex(rocke_lower_t* L, double x)
 {
     /* LLVM spells a float hex constant as the 64-bit hex of the double value
@@ -734,14 +742,6 @@ const char* rocke_ll_fp32_hex(rocke_lower_t* L, double x)
     return rocke_arena_printf(&L->arena, "0x%016llX", (unsigned long long)bits);
 }
 
-const char* rocke_ll_fp64_hex(rocke_lower_t* L, double x)
-{
-    /* LLVM textual IR spells `double` hex constants as the raw IEEE-754 bits. */
-    uint64_t bits;
-    memcpy(&bits, &x, sizeof bits);
-    return rocke_arena_printf(&L->arena, "0x%016llX", (unsigned long long)bits);
-}
-
 const char* rocke_ll_fp16_hex(rocke_lower_t* L, double x)
 {
     /* LLVM IR: half 0xH<4 hex>. Convert double -> IEEE-754 binary16 (round to
@@ -892,17 +892,17 @@ const char* rocke_ll_operand(rocke_lower_t* L, const rocke_value_t* v)
         {
             ity = "i32";
         }
-        if(strcmp(ity, "f32") == 0)
+        if(strcmp(ity, "f64") == 0)
         {
             double fv = 0.0;
             rocke_attr_get_float(&op->attrs, "value", &fv);
-            return rocke_ll_fp32_hex(L, fv);
+            return rocke_ll_fp64_hex(L, fv);
         }
-        if(strcmp(ity, "f64") == 0)
+        if(strcmp(ity, "f32") == 0)
         {
             double fv = 0.0;
             rocke_attr_get_float(&op->attrs, "value", &fv);
-            return rocke_ll_fp64_hex(L, fv);
+            return rocke_ll_fp32_hex(L, fv);
         }
         if(strcmp(ity, "f16") == 0)
         {
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/data.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/data.cpp
index 235b1ad66c14..45f646fabcd6 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/data.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/data.cpp
@@ -150,12 +150,12 @@ const rocke_ll_decl_t ROCKE_LL_INTRINSIC_DECLS[] = {
     {"sqrt.f32", "declare float @llvm.sqrt.f32(float)"},
     {"rsqrt.f32", "declare float @llvm.amdgcn.rsq.f32(float)"},
     {"rcp.f32", "declare float @llvm.amdgcn.rcp.f32(float)"},
-    {"maxnum.f32", "declare float @llvm.maxnum.f32(float, float)"},
     {"maxnum.f64", "declare double @llvm.maxnum.f64(double, double)"},
+    {"maxnum.f32", "declare float @llvm.maxnum.f32(float, float)"},
     {"maxnum.f16", "declare half @llvm.maxnum.f16(half, half)"},
     {"maxnum.bf16", "declare bfloat @llvm.maxnum.bf16(bfloat, bfloat)"},
-    {"minnum.f32", "declare float @llvm.minnum.f32(float, float)"},
     {"minnum.f64", "declare double @llvm.minnum.f64(double, double)"},
+    {"minnum.f32", "declare float @llvm.minnum.f32(float, float)"},
     {"minnum.f16", "declare half @llvm.minnum.f16(half, half)"},
     {"minnum.bf16", "declare bfloat @llvm.minnum.bf16(bfloat, bfloat)"},
     {"fabs.f32", "declare float @llvm.fabs.f32(float)"},
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_llvm_internal.h b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_llvm_internal.h
index 149282e3579d..fd4899864e79 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_llvm_internal.h
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_llvm_internal.h
@@ -565,10 +565,10 @@ bool rocke_ll_is_constant(const rocke_value_t* v);
  * + returns 0 if not a constant. */
 int64_t rocke_ll_eval_constant(rocke_lower_t* L, const rocke_value_t* v);
 
-/* fp32 / fp64 / fp16 LLVM hex constant spellings (Python _fp32_hex / _fp64_hex /
+/* fp64 / fp32 / fp16 LLVM hex constant spellings (Python _fp64_hex / _fp32_hex /
  * _fp16_hex). Arena-owned. */
-const char* rocke_ll_fp32_hex(rocke_lower_t* L, double x);
 const char* rocke_ll_fp64_hex(rocke_lower_t* L, double x);
+const char* rocke_ll_fp32_hex(rocke_lower_t* L, double x);
 const char* rocke_ll_fp16_hex(rocke_lower_t* L, double x);
 
 /* Escape a string for an LLVM asm/string literal (Python
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/ir_model.md b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/ir_model.md
index 3b587f6c6033..07a243a0d9f7 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/ir_model.md
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/ir_model.md
@@ -100,7 +100,7 @@ zero_vec(elem, n)          # f32 / f16 / bf16
 ```text
 add, sub, mul, div, mod                    # integer arithmetic
 fadd, fsub, fmul, fdiv, fneg               # f32 arithmetic
-fmax, fmin                                 # llvm.maxnum / llvm.minnum (f32/f64/f16/bf16)
+fmax, fmin                                 # llvm.maxnum / llvm.minnum (f64/f32/f16/bf16)
 fcmp(pred, a, b)                           # pred in {olt,ole,ogt,oge,oeq,one,ord,uno}
 cmp_lt, cmp_le, cmp_gt, cmp_ge, cmp_eq, cmp_ne   # integer compare (-> i1)
 land, lor, lnot                            # bitwise + i1 logic
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py
index 2cc4f2d7a6ab..4e6b5af40c63 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py
@@ -416,12 +416,12 @@ def _resolve_llvm_flavor() -> str:
     "sqrt.f32": "declare float @llvm.sqrt.f32(float)",
     "rsqrt.f32": "declare float @llvm.amdgcn.rsq.f32(float)",
     "rcp.f32": "declare float @llvm.amdgcn.rcp.f32(float)",
-    "maxnum.f32": "declare float @llvm.maxnum.f32(float, float)",
     "maxnum.f64": "declare double @llvm.maxnum.f64(double, double)",
+    "maxnum.f32": "declare float @llvm.maxnum.f32(float, float)",
     "maxnum.f16": "declare half @llvm.maxnum.f16(half, half)",
     "maxnum.bf16": "declare bfloat @llvm.maxnum.bf16(bfloat, bfloat)",
-    "minnum.f32": "declare float @llvm.minnum.f32(float, float)",
     "minnum.f64": "declare double @llvm.minnum.f64(double, double)",
+    "minnum.f32": "declare float @llvm.minnum.f32(float, float)",
     "minnum.f16": "declare half @llvm.minnum.f16(half, half)",
     "minnum.bf16": "declare bfloat @llvm.minnum.bf16(bfloat, bfloat)",
     "fabs.f32": "declare float @llvm.fabs.f32(float)",
@@ -1565,10 +1565,10 @@ def _operand(self, v: Value) -> str:
         if op.name == "arith.constant":
             ity = op.attrs.get("ity", "i32")
             val = op.attrs["value"]
-            if ity == "f32":
-                return _fp32_hex(val)
             if ity == "f64":
                 return _fp64_hex(val)
+            if ity == "f32":
+                return _fp32_hex(val)
             if ity == "f16":
                 return _fp16_hex(val)
             return str(int(val))
@@ -2248,14 +2248,14 @@ def _op_arith_fmax(self, op: Op) -> None:
         # have their own maxnum intrinsic. The previous f32-only path
         # silently mis-typed half operands as float and broke comgr.
         llvm_ty = {
-            "f32": "float",
             "f64": "double",
+            "f32": "float",
             "f16": "half",
             "bf16": "bfloat",
         }.get(ty_name)
         intrin_key = {
-            "f32": "maxnum.f32",
             "f64": "maxnum.f64",
+            "f32": "maxnum.f32",
             "f16": "maxnum.f16",
             "bf16": "maxnum.bf16",
         }.get(ty_name)
@@ -2271,14 +2271,14 @@ def _op_arith_fmin(self, op: Op) -> None:
         a, b = op.operands
         ty_name = a.type.name
         llvm_ty = {
-            "f32": "float",
             "f64": "double",
+            "f32": "float",
             "f16": "half",
             "bf16": "bfloat",
         }.get(ty_name)
         intrin_key = {
-            "f32": "minnum.f32",
             "f64": "minnum.f64",
+            "f32": "minnum.f32",
             "f16": "minnum.f16",
             "bf16": "minnum.bf16",
         }.get(ty_name)
@@ -6412,6 +6412,14 @@ def _encode_waitcnt_gfx11(vmcnt: int, expcnt: int, lgkmcnt: int) -> int:
     return (ec_b & 0x7) | ((lk_b & 0x3F) << 4) | ((vm_b & 0x3F) << 10)
 
 
+def _fp64_hex(x: float) -> str:
+    # LLVM textual IR spells `double` hex constants as the raw IEEE-754 bits.
+    import struct
+
+    bits = struct.unpack("<Q", struct.pack("<d", float(x)))[0]
+    return f"0x{bits:016X}"
+
+
 def _fp32_hex(x: float) -> str:
     import struct
 
@@ -6424,14 +6432,6 @@ def _fp32_hex(x: float) -> str:
     return f"0x{bits:016X}"
 
 
-def _fp64_hex(x: float) -> str:
-    # LLVM textual IR spells `double` hex constants as the raw IEEE-754 bits.
-    import struct
-
-    bits = struct.unpack("<Q", struct.pack("<d", float(x)))[0]
-    return f"0x{bits:016X}"
-
-
 def _fp16_hex(x: float) -> str:
     # LLVM IR accepts fp16 constants as `half 0xH<4 hex digits>`. We use
     # a numerically-correct rounding via the struct module.
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/distribution.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/distribution.py
index da07477d774d..dbd4445e7633 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/distribution.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/distribution.py
@@ -979,10 +979,10 @@ def store_tile(
     if traits is None:
         traits = make_load_store_traits(distribution)
     quant_dtypes = ("i8", "fp8e4m3", "bf8e5m2")
-    if window.dtype.name not in ("f16", "bf16", "f32", "f64", "i32") + quant_dtypes:
+    if window.dtype.name not in ("f16", "bf16", "f32", "i32", "f64") + quant_dtypes:
         raise NotImplementedError(
             f"store_tile dtype {window.dtype.name} not wired "
-            "(f16/bf16/f32/f64/i32/i8/fp8e4m3/bf8e5m2)"
+            "(f16/bf16/f32/i32/f64/i8/fp8e4m3/bf8e5m2)"
         )
     for y_base in traits.iterate_accesses():
         x_coords = distribution.calculate_x(

From 648ec07398395e1d1ca8d4696ad3de77945c46a7 Mon Sep 17 00:00:00 2001
From: Minseob Shin <minseob.shin11@gmail.com>
Date: Fri, 9 Oct 2026 20:12:26 +0000
Subject: [PATCH 4/7] Reject unsupported f64 ops in HIP, fix f64 vector
 add/sub/mul, extend f64 test coverage

---
 .../platform/cpp/core/lower_hip/core.cpp      |  14 ++
 .../cpp/core/lower_hip/lower_hip_arith.cpp    |  12 ++
 .../cpp/core/lower_hip/lower_hip_mma.cpp      |   2 +
 .../platform/cpp/core/lower_llvm/control.cpp  |   3 +-
 .../cpp/include/rocke/lower_hip_internal.h    |   4 +
 .../platform/dsl_docs/ir_lowering/ir_model.md |   4 +-
 .../dsl_docs/reference/dtype_support.md       |  39 ++--
 .../dsl_docs/reference/op_vocabulary.md       |   6 +-
 .../platform/python/rocke/core/lower_hip.py   |  14 ++
 .../platform/python/rocke/core/lower_llvm.py  |  12 +-
 .../examples/common/hip_lowering_parity.py    |  28 +++
 .../rocke/platform/tests/CMakeLists.txt       |  16 ++
 .../tests/core/test_amdgpu_lowering_ops.py    |  70 ++++++-
 .../platform/tests/core/test_f64_lowering.cpp | 178 ++++++++++++++++++
 .../tests/instances/differential/numeric.py   |   2 +
 .../rocke/platform/tests/test_rocke.py        |  11 --
 16 files changed, 383 insertions(+), 32 deletions(-)
 create mode 100644 dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_f64_lowering.cpp

diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/core.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/core.cpp
index 2e864d78144e..4f763b49b6d9 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/core.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/core.cpp
@@ -416,6 +416,20 @@ const char* rocke_h_vec_prefix_checked(rocke_h_lowerer_t* lw,
     return rocke_h_vec_prefix(ir_scalar_name, full_map);
 }
 
+/* Python _reject_f64(op, tname):
+ *   if tname == "f64":
+ *       raise NotImplementedError(f"{op.name} does not support f64") */
+void rocke_h_reject_f64(rocke_h_lowerer_t* lw, const rocke_op_t* op, const char* tname)
+{
+    if(tname && strcmp(tname, "f64") == 0)
+    {
+        rocke_h_fail(lw,
+                     ROCKE_ERR_NOTIMPL,
+                     "%s does not support f64",
+                     op->name ? op->name : rocke_opcode_name(op->opcode));
+    }
+}
+
 /* Python _type_to_hip(t). Returns an arena-owned string; throws on an unmappable
  * HIP type (KeyError parity), including logical types without direct HIP lowering. */
 const char* rocke_h_type_to_hip(rocke_h_lowerer_t* lw, const rocke_type_t* t)
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_arith.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_arith.cpp
index bf96281ef86a..6afb2e8357aa 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_arith.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_arith.cpp
@@ -517,6 +517,7 @@ static rocke_status_t rocke_h_op_arith_fneg(rocke_h_lowerer_t* lw, const rocke_o
 /* def _op_arith_fabs(self, op):
  *     (v,) = op.operands
  *     ty = _type_to_hip(op.result.type)
+ *     _reject_f64(op, op.result.type.name)
  *     helper = {"f32":"fabsf","f16":"__builtin_fabsf","bf16":"__builtin_fabsf"}
  *              .get(op.result.type.name, "fabsf")
  *     self._emit(f"{ty} {_name(op.result)} = ({ty}){helper}((float){_name(v)});") */
@@ -526,6 +527,7 @@ static rocke_status_t rocke_h_op_arith_fabs(rocke_h_lowerer_t* lw, const rocke_o
     const rocke_value_t* r = h_res(op);
     const char* ty = rocke_h_type_to_hip(lw, r->type);
     const char* tname = r->type->name;
+    rocke_h_reject_f64(lw, op, tname);
     const char* helper = "fabsf";
     if(tname && (strcmp(tname, "f16") == 0 || strcmp(tname, "bf16") == 0))
     {
@@ -539,6 +541,7 @@ static rocke_status_t rocke_h_op_arith_fabs(rocke_h_lowerer_t* lw, const rocke_o
 /* def _op_arith_fma(self, op):
  *     a, b, c = op.operands
  *     ty = _type_to_hip(op.result.type)
+ *     _reject_f64(op, op.result.type.name)
  *     self._emit(f"{ty} {_name(op.result)} = ({ty})fmaf("
  *                f"(float){_name(a)}, (float){_name(b)}, (float){_name(c)});") */
 static rocke_status_t rocke_h_op_arith_fma(rocke_h_lowerer_t* lw, const rocke_op_t* op)
@@ -548,6 +551,7 @@ static rocke_status_t rocke_h_op_arith_fma(rocke_h_lowerer_t* lw, const rocke_op
     const rocke_value_t* c = op->operands[2];
     const rocke_value_t* r = h_res(op);
     const char* ty = rocke_h_type_to_hip(lw, r->type);
+    rocke_h_reject_f64(lw, op, r->type->name);
     rocke_h_emitf(lw,
                   "%s %s = (%s)fmaf((float)%s, (float)%s, (float)%s);",
                   ty,
@@ -562,6 +566,7 @@ static rocke_status_t rocke_h_op_arith_fma(rocke_h_lowerer_t* lw, const rocke_op
 /* def _op_arith_fmax3(self, op):
  *     a, b, c = op.operands
  *     ty = _type_to_hip(op.result.type)
+ *     _reject_f64(op, op.result.type.name)
  *     self._emit(f"{ty} {_name(op.result)} = "
  *                f"(({_name(b)} > {_name(c)}) ? {_name(b)} : {_name(c)});")
  *     self._emit(f"{_name(op.result)} = "
@@ -573,6 +578,7 @@ static rocke_status_t rocke_h_op_arith_fmax3(rocke_h_lowerer_t* lw, const rocke_
     const rocke_value_t* c = op->operands[2];
     const rocke_value_t* r = h_res(op);
     const char* ty = rocke_h_type_to_hip(lw, r->type);
+    rocke_h_reject_f64(lw, op, r->type->name);
     const char* rn = rocke_h_name(lw, r);
     const char* an = rocke_h_name(lw, a);
     const char* bn = rocke_h_name(lw, b);
@@ -585,6 +591,7 @@ static rocke_status_t rocke_h_op_arith_fmax3(rocke_h_lowerer_t* lw, const rocke_
 /* def _op_arith_fmin3(self, op):
  *     a, b, c = op.operands
  *     ty = _type_to_hip(op.result.type)
+ *     _reject_f64(op, op.result.type.name)
  *     self._emit(f"{ty} {_name(op.result)} = "
  *                f"(({_name(b)} < {_name(c)}) ? {_name(b)} : {_name(c)});")
  *     self._emit(f"{_name(op.result)} = "
@@ -596,6 +603,7 @@ static rocke_status_t rocke_h_op_arith_fmin3(rocke_h_lowerer_t* lw, const rocke_
     const rocke_value_t* c = op->operands[2];
     const rocke_value_t* r = h_res(op);
     const char* ty = rocke_h_type_to_hip(lw, r->type);
+    rocke_h_reject_f64(lw, op, r->type->name);
     const char* rn = rocke_h_name(lw, r);
     const char* an = rocke_h_name(lw, a);
     const char* bn = rocke_h_name(lw, b);
@@ -737,6 +745,7 @@ static rocke_status_t rocke_h_op_arith_fcmp(rocke_h_lowerer_t* lw, const rocke_o
  *       (v,) = op.operands
  *       tname = op.result.type.name
  *       cpp_t = _type_to_hip(op.result.type)
+ *       _reject_f64(op, tname)
  *       if tname == "f32":
  *           self._emit(f"{cpp_t} {_name(op.result)} = {fn_f32}({_name(v)});")
  *       else:
@@ -750,6 +759,7 @@ static void h_math1(rocke_h_lowerer_t* lw, const rocke_op_t* op, const char* fn_
     const char* cpp_t = rocke_h_type_to_hip(lw, r->type);
     const char* rn = rocke_h_name(lw, r);
     const char* vn = rocke_h_name(lw, v);
+    rocke_h_reject_f64(lw, op, tname);
     if(tname && strcmp(tname, "f32") == 0)
     {
         rocke_h_emitf(lw, "%s %s = %s(%s);", cpp_t, rn, fn_f32, vn);
@@ -763,6 +773,7 @@ static void h_math1(rocke_h_lowerer_t* lw, const rocke_op_t* op, const char* fn_
 /* An amdgcn-builtin reciprocal-style math op (rcp / rcp_fast / rsqrt). Same
  * promote-compute-demote shape as _math1 but the f32 path drops the (cpp_t)
  * cast since the builtin already returns float:
+ *   _reject_f64(op, tname)
  *   if tname == "f32":
  *       self._emit(f"{cpp_t} {_name} = {builtin}({_name(v)});")
  *   else:
@@ -775,6 +786,7 @@ static void h_amdgcn_unary(rocke_h_lowerer_t* lw, const rocke_op_t* op, const ch
     const char* cpp_t = rocke_h_type_to_hip(lw, r->type);
     const char* rn = rocke_h_name(lw, r);
     const char* vn = rocke_h_name(lw, v);
+    rocke_h_reject_f64(lw, op, tname);
     if(tname && strcmp(tname, "f32") == 0)
     {
         rocke_h_emitf(lw, "%s %s = %s(%s);", cpp_t, rn, builtin, vn);
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mma.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mma.cpp
index 6abf236d5547..5cc77c73348c 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mma.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mma.cpp
@@ -1069,6 +1069,8 @@ static rocke_status_t rocke_h_op_vector_fma(rocke_h_lowerer_t* lw, const rocke_o
     const rocke_value_t* c = op->operands[2];
     const rocke_value_t* r = h_res(op);
     const char* res_t = rocke_h_type_to_hip(lw, r->type);
+    rocke_h_reject_f64(
+        lw, op, r->type->kind == ROCKE_TYPE_VECTOR ? r->type->elem->name : r->type->name);
     const char* nice = rocke_h_name(lw, r);
     const char* an = rocke_h_name(lw, a);
     const char* bn = rocke_h_name(lw, b);
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/control.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/control.cpp
index 1dd8228e9465..4706b369f67f 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/control.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/control.cpp
@@ -60,7 +60,8 @@ static const char* ll_vec_elem_name(const rocke_type_t* t)
 static bool ll_is_fp_elem(const char* name)
 {
     return name
-           && (strcmp(name, "f16") == 0 || strcmp(name, "bf16") == 0 || strcmp(name, "f32") == 0);
+           && (strcmp(name, "f16") == 0 || strcmp(name, "bf16") == 0 || strcmp(name, "f32") == 0
+               || strcmp(name, "f64") == 0);
 }
 
 /* ======================================================================== */
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_hip_internal.h b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_hip_internal.h
index ae2283f8c2c1..3192eabf6796 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_hip_internal.h
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_hip_internal.h
@@ -172,6 +172,10 @@ const char* rocke_h_vec_prefix_checked(rocke_h_lowerer_t* lw,
                                        bool full_map,
                                        const char* op_desc);
 
+/* Python _reject_f64(op, tname): raises ROCKE_ERR_NOTIMPL "<op> does not support
+ * f64" when `tname` is "f64", for an op that has no f64 lowering. */
+void rocke_h_reject_f64(rocke_h_lowerer_t* lw, const rocke_op_t* op, const char* tname);
+
 /* --------------------------------------------------------- literals / encode */
 
 /* Python _f32_literal(val): C++ float-literal text for a double, special-casing
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/ir_model.md b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/ir_model.md
index 07a243a0d9f7..570729eb16f3 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/ir_model.md
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/ir_model.md
@@ -99,7 +99,7 @@ zero_vec(elem, n)          # f32 / f16 / bf16
 
 ```text
 add, sub, mul, div, mod                    # integer arithmetic
-fadd, fsub, fmul, fdiv, fneg               # f32 arithmetic
+fadd, fsub, fmul, fdiv, fneg               # f64/f32/f16/bf16 arithmetic
 fmax, fmin                                 # llvm.maxnum / llvm.minnum (f64/f32/f16/bf16)
 fcmp(pred, a, b)                           # pred in {olt,ole,ogt,oge,oeq,one,ord,uno}
 cmp_lt, cmp_le, cmp_gt, cmp_ge, cmp_eq, cmp_ne   # integer compare (-> i1)
@@ -126,7 +126,7 @@ cvt_f32_to_i8_sat(v)     # round + saturate to i8
 
 ### Math intrinsics
 
-LLVM target intrinsics; lowering is f32-centric. `exp2`, `sqrt`, `rsqrt` and `tanh` accept only f32 and reject f64 inputs; `rcp` is a plain `fdiv` and accepts f64.
+LLVM target intrinsics; lowering is f32-centric. `exp2`, `exp2_fast`, `log2`, `sqrt`, `rsqrt`, `rcp_fast`, `tanh`, `fabs`, `fma`, `fmax3`, `fmin3` and `vector_fma` have no f64 lowering and reject f64 inputs in both the LLVM and HIP lowerings; `rcp` is a plain `fdiv` and accepts f64.
 
 ```text
 exp2(v)        # llvm.exp2.f32
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/dtype_support.md b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/dtype_support.md
index 1e5804d50b11..bb0167165e15 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/dtype_support.md
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/dtype_support.md
@@ -27,14 +27,18 @@ estimates, shipping order, and acceptance checklists are maintained separately.
 logical dtype and native gfx942 atoms added after the original basis snapshot.
 Other scores and source line numbers retain their stated historical basis.
 
+**FP64 coverage update.** The f64 entries in §1, §2.5, and §4C reflect the IR
+type, arithmetic, memory, and elementwise/reduce2d support added after the
+original basis snapshot. The §2.4 f64 row retains its stated historical basis.
+
 ---
 
 ## 1. The canonical dtype set
 
 **Current clarification (2026-10-07).** The original basis snapshot listed
 ten directly lowerable scalar representations. The table below also includes
-the subsequently added logical TF32 type with its I32 carrier. The registry
-recognizes `fp4e2m1`, `fp6e2m3`, `fp6e3m2`, `e8m0`, and `e5m3` as logical `Type`
+the subsequently added logical TF32 type with its I32 carrier and the F64
+type. The registry recognizes `fp4e2m1`, `fp6e2m3`, `fp6e3m2`, `e8m0`, and `e5m3` as logical `Type`
 identities preserved by serialization; their packed-data paths use `i8` storage
 and integer carrier registers. See [dtype storage and packing](../architecture/dtype_storage_packing.md).
 Logical recognition does not provide generic scalar arithmetic or conversion.
@@ -59,6 +63,7 @@ The definitions in both engines are:
 | `bf16` | `bfloat` | 16 | float |
 | `f32` | `float` | 32 | float |
 | `tf32` | `i32` (carrier) | 32 | logical float; native gfx942 XF32 input |
+| `f64` | `double` | 64 | float |
 | `fp8e4m3` | `i8` (storage) | 8 | float (OCP E4M3) |
 | `bf8e5m2` | `i8` (storage) | 8 | float (OCP E5M2) |
 
@@ -117,7 +122,7 @@ equivalent; `USE_TF32` output is not a bit-accuracy oracle for BF16 compute.
 The matrix below therefore labels this possible mode `bf16-mode`. It is separate
 from logical TF32; no production selector maps TF32 requests to BF16 compute.
 
-**Types that are entirely absent from the vocabulary:** `f64`, unsigned
+**Types that are entirely absent from the vocabulary:** unsigned
 `u8/u16/u32/u64`, and a distinct `bool` (predicates reuse `i1`). No engine
 references them.
 
@@ -196,6 +201,8 @@ Scores 0–3 per criterion; `—` = N/A (excluded from the denominator). Rows ar
 grouped by type family, not ranked — see "What the scores are not" in §2.3.
 TF32 was added after this scored snapshot; its current layer boundaries are
 described in §1 and §4D rather than assigned a retrospective percentage here.
+The `f64` row records the snapshot before F64 was added; its current layer
+boundaries are described in §2.5.
 
 Gap categories: **[C]** conversion gap · **[P]** partial/plumbing gap ·
 **[M]** missing entirely.
@@ -338,11 +345,21 @@ mode. The missing selector and precision contract are described in §2.5.
   `cast_f32_to(v, BF16)` (C3=2), and bf16 MFMA/WMMA atoms (C2=3).
   **C6=1** records the need for a caller-visible BF16 precision contract when
   the mode is exposed; see [§1](#1-the-canonical-dtype-set).
-- **`f64` / `u16-64` — 0 at every layer.** Absent from the hipDNN enum, rejected
-  by the MIOpen provider, and not requested by any family. For `f64` specifically
-  `V_MFMA_F64_*` is present on gfx90a/gfx942/gfx950, where f64 is unwired
-  software. gfx1250 has no F64 matrix instruction: LLVM gates
-  `V_WMMA_F64_16X16X4_F64` on `gfx1251-gemm-insts` (§4A).
+- **`u16-64` — 0 at every layer.** Absent from the hipDNN enum, rejected by the
+  MIOpen provider, and not requested by any family.
+- **`f64` — IR, memory, and elementwise/reduce2d only.** The IR type,
+  `const_f64`, f64 `fadd`/`fsub`/`fmul`/`fdiv`/`fneg`/`fmax`/`fmin`/`fcmp`/
+  `select` and `vector_add`/`vector_sub`/`vector_mul`, global/LDS memory up to
+  2-element vectors, and HIP lowering are present. Elementwise
+  (non-transcendental ops) and reduce2d accept f64 with `vec=2`
+  ([SUPPORT_MATRIX](../../python/rocke/instances/SUPPORT_MATRIX.md)).
+  `exp2`, `exp2_fast`, `log2`, `sqrt`, `rsqrt`, `rcp_fast`, `tanh`, `fabs`,
+  `fma`, `fmax3`, `fmin3`, and `vector_fma` reject f64 in both lowerings; there
+  are no conversions to or from f64, MMA atoms, or atomics. f64 is in the
+  hipDNN enum (`DOUBLE`) but rejected by the MIOpen provider.
+  `V_MFMA_F64_*` is present on gfx90a/gfx942/gfx950 but undeclared (§4C.6).
+  gfx1250 has no F64 matrix instruction: LLVM gates `V_WMMA_F64_16X16X4_F64` on
+  `gfx1251-gemm-insts` (§4A).
 
 ### 2.6 Source anchors for the scores
 
@@ -462,7 +479,7 @@ The original basis snapshot, with the gfx942 XF32 addition noted below:
 ### 4C. Deltas — silicon capability vs. rocKE's declared atoms
 
 The following gaps compare §4A with §4B. The first five concern atom declarations
-only; f64 (4C.6) also lacks an IR type. The gfx942 XF32 declaration gap is now
+only, as does f64 (4C.6). The gfx942 XF32 declaration gap is now
 closed; §4D describes its logical dtype and remaining integration boundaries.
 
 #### 4C.1 gfx1250 — fp4/fp6/bf6 block-scaled atoms are omitted (deliberately)
@@ -498,8 +515,8 @@ catalog declares no sparse atom on any arch, and no family requests one. No dtyp
 in §2.4 is blocked on it.
 
 #### 4C.6 f64 matrix is in silicon but undeclared
-`V_MFMA_F64_*` is present on gfx90a/gfx942/gfx950. `f64` is not an IR type
-at all (§1), so those targets have a vocabulary gap as well as missing atoms.
+`V_MFMA_F64_*` is present on gfx90a/gfx942/gfx950. `f64` is an IR type
+(§1), but no f64 MMA atom is declared on those targets.
 gfx1250 has no native F64 matrix instruction and is excluded from this gap.
 
 ### 4D. XF32 / TF32 across the arch set
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md
index 3341a2835204..b90b994e3dff 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md
@@ -19,8 +19,8 @@ Complete reference of operations recognized by `core/ir.py` and lowered to AMDGP
 | `arith.fmul`           | `fmul`                            | `fmul`                                 |
 | `arith.fdiv`           | `fdiv`                            | `fdiv`                                 |
 | `arith.fneg`           | `fneg`                            | `fneg`                                 |
-| `arith.fmax`           | `fmax`                            | `llvm.maxnum.{f32,f16,bf16}`           |
-| `arith.fmin`           | `fmin`                            | `llvm.minnum.{f32,f16,bf16}`           |
+| `arith.fmax`           | `fmax`                            | `llvm.maxnum.{f64,f32,f16,bf16}`       |
+| `arith.fmin`           | `fmin`                            | `llvm.minnum.{f64,f32,f16,bf16}`       |
 | `arith.cmp`            | `cmp_lt`, `cmp_le`, `cmp_gt`, `cmp_ge`, `cmp_eq`, `cmp_ne` | `icmp {slt,sle,sgt,sge,eq,ne}` |
 | `arith.fcmp`           | `fcmp(pred, ...)`                 | `fcmp {olt,ole,ogt,oge,oeq,one,ord,uno}` |
 | `arith.and`            | `land`                            | `and`                                  |
@@ -51,7 +51,7 @@ Complete reference of operations recognized by `core/ir.py` and lowered to AMDGP
 | `math.tanh`    | `tanh`       | piecewise f32 expansion: OCML polynomial for `abs(x) < 0.625`, otherwise `exp2` + reciprocal; bitwise sign restoration |
 | `math.rcp`     | `rcp`        | `1.0 / v` (hardware reciprocal) |
 
-`exp2`, `sqrt`, `rsqrt` and `tanh` are f32-only and reject f64 inputs; `rcp` accepts f64.
+`exp2`, `exp2_fast`, `log2`, `sqrt`, `rsqrt`, `rcp_fast`, `tanh`, `fabs`, `fma`, `fmax3`, `fmin3` and `vector_fma` have no f64 lowering and reject f64 inputs in both the LLVM and HIP lowerings; `rcp` accepts f64.
 
 `clamp_f32(v, lo, hi)` is `fmin(hi, fmax(lo, v))` — folds to `v_med3_f32`.
 
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py
index a3a3a10028cb..3b2dc303e470 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py
@@ -223,6 +223,12 @@ def _vec_prefix(elem_name: str, op_desc: str) -> str:
         ) from None
 
 
+def _reject_f64(op: Op, tname: str) -> None:
+    """Reject an f64 operand for an op that has no f64 lowering."""
+    if tname == "f64":
+        raise NotImplementedError(f"{op.name} does not support f64")
+
+
 def _f32_literal(val: float) -> str:
     """Format a Python float for C++ float literal context.
 
@@ -1324,6 +1330,7 @@ def _op_arith_fneg(self, op: Op) -> None:
     def _op_arith_fabs(self, op: Op) -> None:
         (v,) = op.operands
         ty = _type_to_hip(op.result.type)
+        _reject_f64(op, op.result.type.name)
         # `__builtin_fabsf` lowers to a free `abs` modifier on AMDGPU
         # (exactly what the LLVM `llvm.fabs.f32` intrinsic produces);
         # for f16/bf16 the matching builtin yields the same modifier.
@@ -1337,6 +1344,7 @@ def _op_arith_fabs(self, op: Op) -> None:
     def _op_arith_fma(self, op: Op) -> None:
         a, b, c = op.operands
         ty = _type_to_hip(op.result.type)
+        _reject_f64(op, op.result.type.name)
         # `fmaf(a, b, c)` lowers to `v_fma_f32` on AMDGPU for f32; for
         # f16/bf16 we promote to f32, fma, and demote — semantically
         # identical to the IR-level `arith.fma` and matches the
@@ -1349,6 +1357,7 @@ def _op_arith_fma(self, op: Op) -> None:
     def _op_arith_fmax3(self, op: Op) -> None:
         a, b, c = op.operands
         ty = _type_to_hip(op.result.type)
+        _reject_f64(op, op.result.type.name)
         # Two ternary ops in one statement; on AMDGPU the chain folds
         # into a single `v_max3_f32`.
         self._emit(
@@ -1363,6 +1372,7 @@ def _op_arith_fmax3(self, op: Op) -> None:
     def _op_arith_fmin3(self, op: Op) -> None:
         a, b, c = op.operands
         ty = _type_to_hip(op.result.type)
+        _reject_f64(op, op.result.type.name)
         self._emit(
             f"{ty} {_name(op.result)} = "
             f"(({_name(b)} < {_name(c)}) ? {_name(b)} : {_name(c)});"
@@ -1441,6 +1451,7 @@ def _math1(
         (v,) = op.operands
         tname = op.result.type.name
         cpp_t = _type_to_hip(op.result.type)
+        _reject_f64(op, tname)
         # ``__builtin_amdgcn_*`` only takes float. For non-f32 types,
         # round-trip via float so the math runs at f32 precision.
         if tname == "f32":
@@ -1485,6 +1496,7 @@ def _op_math_rcp_fast(self, op: Op) -> None:
         (v,) = op.operands
         tname = op.result.type.name
         cpp_t = _type_to_hip(op.result.type)
+        _reject_f64(op, tname)
         if tname == "f32":
             self._emit(
                 f"{cpp_t} {_name(op.result)} = __builtin_amdgcn_rcpf({_name(v)});"
@@ -1503,6 +1515,7 @@ def _op_math_rsqrt(self, op: Op) -> None:
         (v,) = op.operands
         tname = op.result.type.name
         cpp_t = _type_to_hip(op.result.type)
+        _reject_f64(op, tname)
         if tname == "f32":
             self._emit(
                 f"{cpp_t} {_name(op.result)} = __builtin_amdgcn_rsqf({_name(v)});"
@@ -2477,6 +2490,7 @@ def _op_vector_fma(self, op: Op) -> None:
         a, bb, cc = op.operands
         n = op.result.type.count if isinstance(op.result.type, VectorType) else 1
         res_t = _type_to_hip(op.result.type)
+        _reject_f64(op, op.result.type.elem.name)
         nice = _name(op.result)
         self._emit(f"{res_t} {nice};")
         for i in range(n):
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py
index 4e6b5af40c63..0d84d8c84a69 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py
@@ -5655,15 +5655,21 @@ def _vector_binop(self, op: Op, llvm_op: str) -> None:
 
     def _op_vector_add(self, op: Op) -> None:
         elem = op.result.type.elem.name  # type: ignore[attr-defined]
-        self._vector_binop(op, "fadd" if elem in ("f16", "bf16", "f32") else "add")
+        self._vector_binop(
+            op, "fadd" if elem in ("f16", "bf16", "f32", "f64") else "add"
+        )
 
     def _op_vector_mul(self, op: Op) -> None:
         elem = op.result.type.elem.name  # type: ignore[attr-defined]
-        self._vector_binop(op, "fmul" if elem in ("f16", "bf16", "f32") else "mul")
+        self._vector_binop(
+            op, "fmul" if elem in ("f16", "bf16", "f32", "f64") else "mul"
+        )
 
     def _op_vector_sub(self, op: Op) -> None:
         elem = op.result.type.elem.name  # type: ignore[attr-defined]
-        self._vector_binop(op, "fsub" if elem in ("f16", "bf16", "f32") else "sub")
+        self._vector_binop(
+            op, "fsub" if elem in ("f16", "bf16", "f32", "f64") else "sub"
+        )
 
     def _op_vector_and(self, op: Op) -> None:
         self._vector_binop(op, "and")
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/examples/common/hip_lowering_parity.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/examples/common/hip_lowering_parity.py
index 0f5e0c83d33f..afdeee46ae05 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/examples/common/hip_lowering_parity.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/examples/common/hip_lowering_parity.py
@@ -531,6 +531,17 @@ def make_cases(*, arch: str = "gfx950") -> List[Case]:
                 )
             )
 
+    for op in ("copy", "neg", "abs", "relu", "add", "sub", "mul", "max", "min"):
+        cases.append(
+            Case(
+                f"elementwise.f64.{op}",
+                "small",
+                lambda op=op: build_elementwise(
+                    ElementwiseSpec(op=op, dtype="f64", vec=2)
+                ),
+            )
+        )
+
     for dtype in ("f16", "bf16"):
         cases.extend(
             [
@@ -601,6 +612,23 @@ def make_cases(*, arch: str = "gfx950") -> List[Case]:
             ),
         )
     )
+    for op in ("sum", "max", "min", "mean", "prod"):
+        cases.append(
+            Case(
+                f"reduce.f64.{op}",
+                "small",
+                lambda op=op: build_reduce2d(
+                    Reduce2DSpec(
+                        n_per_block=2048,
+                        op=op,
+                        dtype="f64",
+                        block_size=256,
+                        vec=2,
+                        wave_size=red_wave,
+                    )
+                ),
+            )
+        )
 
     for op in ("avg", "sum"):
         cases.append(
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/CMakeLists.txt b/dnn-providers/hip-kernel-provider/rocke/platform/tests/CMakeLists.txt
index 246352ff980a..9f0947234c31 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/CMakeLists.txt
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/CMakeLists.txt
@@ -156,6 +156,22 @@ if(BUILD_TESTING)
                  rocke_tanh_lowering)
   endif()
 
+  # f64 op rejection and f64 vector arithmetic in both C++ lowerers.
+  if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/core/test_f64_lowering.cpp")
+    add_executable(rocke_f64_lowering core/test_f64_lowering.cpp)
+    target_link_libraries(rocke_f64_lowering PRIVATE rocke_core)
+    add_test(NAME rocke_f64_lowering COMMAND rocke_f64_lowering)
+    set_tests_properties(rocke_f64_lowering
+                         PROPERTIES LABELS "ckc;unit;host")
+    rocke_cxx_test_installed_name(_rocke_installed_name rocke_f64_lowering)
+    set_target_properties(rocke_f64_lowering
+                          PROPERTIES OUTPUT_NAME "${_rocke_installed_name}")
+    install(TARGETS rocke_f64_lowering
+            RUNTIME DESTINATION ${ROCKE_CXX_TEST_INSTALL_DIR})
+    set_property(GLOBAL APPEND PROPERTY ROCKE_INSTALL_TEST_TARGETS
+                 rocke_f64_lowering)
+  endif()
+
   # Holds the ISA-layer and lower_llvm-layer flavor->datalayout tables together.
   if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/core/flavor_datalayout_ssot.cpp")
     add_executable(rocke_flavor_datalayout_ssot core/flavor_datalayout_ssot.cpp)
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_amdgpu_lowering_ops.py b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_amdgpu_lowering_ops.py
index e940a146ae30..48b6f1a8b10e 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_amdgpu_lowering_ops.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_amdgpu_lowering_ops.py
@@ -14,6 +14,11 @@
   map silently reinterpreted them as f16), and reject anything unmapped.
 * ``lower_hip`` arch seam -- an explicitly passed arch must resolve in the arch
   catalog; only an omitted arch falls back to the gfx950 baseline.
+* f64 ops without an f64 lowering (transcendental math, ``fabs``, ``fma``,
+  ``fmax3``, ``fmin3``, ``vector_fma``) -- raise in both ``lower_hip`` and
+  ``lower_llvm`` instead of computing through f32.
+* ``lower_llvm`` f64 ``vector_add`` / ``vector_sub`` / ``vector_mul`` -- emit
+  ``fadd`` / ``fsub`` / ``fmul`` on ``<N x double>``, not the integer ops.
 * ``lower_llvm._op_memref_global_atomic_add_pk_bf16`` -- lowers to a generic
   ``atomicrmw fadd <2 x bfloat>`` (the ``llvm.amdgcn.global.atomic.fadd.v2bf16``
   intrinsic does not exist in the shipping ROCm LLVM).
@@ -31,7 +36,7 @@
 
 import unittest
 
-from rocke.core.ir import BF16, I32, IRBuilder, PtrType
+from rocke.core.ir import BF16, F64, I32, IRBuilder, PtrType
 from rocke.core.lower_hip import lower_kernel_to_hip
 from rocke.core.lower_llvm import lower_kernel_to_llvm
 
@@ -160,6 +165,69 @@ def test_omitted_arch_uses_baseline(self):
         )
 
 
+class TestHipF64Rejection(unittest.TestCase):
+    """f64 ops without an f64 lowering fail explicitly in both lowerers."""
+
+    def _kernel(self, op_name):
+        b = IRBuilder(f"{op_name}_f64")
+        out = b.param("out", PtrType(F64, "global"))
+        idx = b.const_i32(0)
+        if op_name == "vector_fma":
+            v = b.global_load_vN(out, idx, F64, 2, align=16)
+            b.global_store_vN(out, idx, b.vector_fma(v, v, v), n=2, align=16)
+        else:
+            x = b.global_load(out, idx, F64, align=8)
+            args = (x, x, x) if op_name in ("fma", "fmax3", "fmin3") else (x,)
+            b.global_store(out, idx, getattr(b, op_name)(*args), align=8)
+        b.ret()
+        return b.kernel
+
+    def test_unsupported_ops_raise(self):
+        for op_name in (
+            "exp2",
+            "log2",
+            "sqrt",
+            "rsqrt",
+            "rcp_fast",
+            "fabs",
+            "fma",
+            "fmax3",
+            "fmin3",
+            "vector_fma",
+        ):
+            with self.subTest(op=op_name):
+                kernel = self._kernel(op_name)
+                with self.assertRaises(NotImplementedError) as cm:
+                    lower_kernel_to_hip(kernel, arch="gfx942")
+                self.assertIn("does not support f64", str(cm.exception))
+                with self.assertRaises(NotImplementedError):
+                    lower_kernel_to_llvm(kernel, arch="gfx942")
+
+    def test_rcp_divides_in_double(self):
+        hip = lower_kernel_to_hip(self._kernel("rcp"), arch="gfx942")
+        self.assertIn("= 1.0 / ", hip)
+        self.assertNotIn("__builtin_amdgcn_rcpf", hip)
+
+
+class TestF64VectorArith(unittest.TestCase):
+    def test_vector_binops_use_float_ops(self):
+        for op_name, llvm_op in (
+            ("vector_add", "fadd"),
+            ("vector_sub", "fsub"),
+            ("vector_mul", "fmul"),
+        ):
+            with self.subTest(op=op_name):
+                b = IRBuilder(f"{op_name}_f64")
+                out = b.param("out", PtrType(F64, "global"))
+                idx = b.const_i32(0)
+                v = b.global_load_vN(out, idx, F64, 2, align=16)
+                r = getattr(b, op_name)(v, v)
+                b.global_store_vN(out, idx, r, n=2, align=16)
+                b.ret()
+                ir = lower_kernel_to_llvm(b.kernel, arch="gfx942")
+                self.assertIn(f"{r.name} = {llvm_op} <2 x double>", ir)
+
+
 class TestDsReadTr16NameHint(unittest.TestCase):
     def test_tr16_result_name_avoids_trunc_collision(self):
         b = IRBuilder("tr16_name")
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_f64_lowering.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_f64_lowering.cpp
new file mode 100644
index 000000000000..c685a7cf37a6
--- /dev/null
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_f64_lowering.cpp
@@ -0,0 +1,178 @@
+// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+/*
+ * test_f64_lowering.cpp -- C++ engine coverage for the f64 op boundary.
+ * Mirrors TestHipF64Rejection / TestF64VectorArith in
+ * tests/core/test_amdgpu_lowering_ops.py.
+ */
+
+#include "rocke/ir.h"
+#include "rocke/lower_hip.h"
+#include "rocke/lower_llvm.h"
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+
+#define CHECK(expr)                                                 \
+    do                                                              \
+    {                                                               \
+        if(!(expr))                                                 \
+        {                                                           \
+            std::fprintf(stderr, "line %d: %s\n", __LINE__, #expr); \
+            return 1;                                               \
+        }                                                           \
+    } while(0)
+
+enum f64_op
+{
+    OP_EXP2,
+    OP_LOG2,
+    OP_SQRT,
+    OP_RSQRT,
+    OP_RCP_FAST,
+    OP_FABS,
+    OP_FMA,
+    OP_FMAX3,
+    OP_FMIN3,
+    OP_VECTOR_FMA,
+    OP_RCP,
+    OP_VECTOR_ADD,
+    OP_VECTOR_SUB,
+    OP_VECTOR_MUL,
+};
+
+/* f64 `out[0] = op(out[0], ...)`; vector ops use a 2-wide load/store. */
+static rocke_value_t* build_f64_kernel(rocke_ir_builder_t* b, f64_op op)
+{
+    rocke_value_t* out = rocke_b_param(b, "out", rocke_ptr_type(b, rocke_f64(), "global"), NULL);
+    rocke_value_t* idx = rocke_b_const_i32(b, 0);
+    rocke_value_t* r = NULL;
+    if(op >= OP_VECTOR_FMA && op != OP_RCP)
+    {
+        rocke_value_t* v = rocke_b_global_load_vN(b, out, idx, rocke_f64(), 2, 16);
+        r = op == OP_VECTOR_FMA   ? rocke_b_vector_fma(b, v, v, v)
+            : op == OP_VECTOR_ADD ? rocke_b_vector_add(b, v, v)
+            : op == OP_VECTOR_SUB ? rocke_b_vector_sub(b, v, v)
+                                  : rocke_b_vector_mul(b, v, v);
+        rocke_b_global_store_vN(b, out, idx, r, 2, 16);
+    }
+    else
+    {
+        rocke_value_t* x = rocke_b_global_load(b, out, idx, rocke_f64(), 8);
+        switch(op)
+        {
+        case OP_EXP2:
+            r = rocke_b_exp2(b, x);
+            break;
+        case OP_LOG2:
+            r = rocke_b_log2(b, x);
+            break;
+        case OP_SQRT:
+            r = rocke_b_sqrt(b, x);
+            break;
+        case OP_RSQRT:
+            r = rocke_b_rsqrt(b, x);
+            break;
+        case OP_RCP_FAST:
+            r = rocke_b_rcp_fast(b, x);
+            break;
+        case OP_FABS:
+            r = rocke_b_fabs(b, x);
+            break;
+        case OP_FMA:
+            r = rocke_b_fma(b, x, x, x);
+            break;
+        case OP_FMAX3:
+            r = rocke_b_fmax3(b, x, x, x);
+            break;
+        case OP_FMIN3:
+            r = rocke_b_fmin3(b, x, x, x);
+            break;
+        default:
+            r = rocke_b_rcp(b, x);
+            break;
+        }
+        rocke_b_global_store(b, out, idx, r, 8);
+    }
+    rocke_b_ret(b);
+    return r;
+}
+
+static int test_unsupported_ops_rejected()
+{
+    for(int op = OP_EXP2; op <= OP_VECTOR_FMA; ++op)
+    {
+        rocke_ir_builder_t b;
+        CHECK(rocke_ir_builder_init(&b, "f64_reject") == ROCKE_OK);
+        build_f64_kernel(&b, static_cast<f64_op>(op));
+        CHECK(rocke_ir_builder_ok(&b));
+
+        rocke_strbuf_t hip;
+        CHECK(rocke_strbuf_init(&hip, 0) == 0);
+        rocke_lower_hip_opts_t opts = {};
+        opts.arch = "gfx942";
+        CHECK(rocke_lower_kernel_to_hip(&b, b.kernel, &opts, &hip) == ROCKE_ERR_NOTIMPL);
+        rocke_strbuf_free(&hip);
+
+        char* llvm_text = NULL;
+        CHECK(rocke_lower_kernel_to_llvm(b.kernel, ROCKE_LLVM_FLAVOR_LLVM20, "gfx942", &llvm_text)
+              == ROCKE_ERR_NOTIMPL);
+        std::free(llvm_text);
+        rocke_ir_builder_free(&b);
+    }
+    return 0;
+}
+
+static int test_rcp_divides_in_double()
+{
+    rocke_ir_builder_t b;
+    CHECK(rocke_ir_builder_init(&b, "f64_rcp") == ROCKE_OK);
+    build_f64_kernel(&b, OP_RCP);
+
+    rocke_strbuf_t hip;
+    CHECK(rocke_strbuf_init(&hip, 0) == 0);
+    rocke_lower_hip_opts_t opts = {};
+    opts.arch = "gfx942";
+    CHECK(rocke_lower_kernel_to_hip(&b, b.kernel, &opts, &hip) == ROCKE_OK);
+    CHECK(std::strstr(hip.data, "= 1.0 / ") != NULL);
+    CHECK(std::strstr(hip.data, "__builtin_amdgcn_rcpf") == NULL);
+    rocke_strbuf_free(&hip);
+    rocke_ir_builder_free(&b);
+    return 0;
+}
+
+static int test_vector_binops_use_float_ops()
+{
+    const struct
+    {
+        f64_op op;
+        const char* llvm_op;
+    } cases[] = {{OP_VECTOR_ADD, "fadd"}, {OP_VECTOR_SUB, "fsub"}, {OP_VECTOR_MUL, "fmul"}};
+    for(const auto& c : cases)
+    {
+        rocke_ir_builder_t b;
+        CHECK(rocke_ir_builder_init(&b, "f64_vector") == ROCKE_OK);
+        rocke_value_t* r = build_f64_kernel(&b, c.op);
+
+        char* llvm_text = NULL;
+        CHECK(rocke_lower_kernel_to_llvm(b.kernel, ROCKE_LLVM_FLAVOR_LLVM20, "gfx942", &llvm_text)
+              == ROCKE_OK);
+        char expected[64];
+        std::snprintf(expected, sizeof(expected), "%s = %s <2 x double>", r->name, c.llvm_op);
+        CHECK(std::strstr(llvm_text, expected) != NULL);
+        std::free(llvm_text);
+        rocke_ir_builder_free(&b);
+    }
+    return 0;
+}
+
+int main()
+{
+    if(test_unsupported_ops_rejected() || test_rcp_divides_in_double()
+       || test_vector_binops_use_float_ops())
+    {
+        std::fprintf(stderr, "f64 lowering contract failed\n");
+        return 1;
+    }
+    return 0;
+}
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/numeric.py b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/numeric.py
index 724a1df05ecc..4bb19ca3f1e1 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/numeric.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/numeric.py
@@ -555,9 +555,11 @@ class ElemCfg:
     ElemCfg("tanh_bf16_128k", 2048 * 64, "tanh", "bf16"),
     # f64: native f64 ops only, vec=2 (256*2 = 512 elems/block).
     ElemCfg("add_f64_128k", 2048 * 64, "add", "f64", vec=2),
+    ElemCfg("sub_f64_128k", 2048 * 64, "sub", "f64", vec=2),
     ElemCfg("mul_f64_128k", 2048 * 64, "mul", "f64", vec=2),
     ElemCfg("min_f64_128k", 2048 * 64, "min", "f64", vec=2),
     ElemCfg("relu_f64_128k", 2048 * 64, "relu", "f64", vec=2),
+    ElemCfg("abs_f64_128k", 2048 * 64, "abs", "f64", vec=2),
 ]
 
 # map spec dtype -> TOL-table key
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py b/dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py
index 2d05bb71611e..712a7fc68988 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py
@@ -1094,17 +1094,6 @@ def test_f64_path_builds(self):
             ll = lower_kernel_to_llvm(build_reduce2d(spec))
             self.assertIn("double", ll)
 
-    def test_f64_max_min_identities_are_infinite(self):
-        from rocke.instances import Reduce2DSpec, build_reduce2d
-
-        for op, inf_bits in (
-            ("max", "0xFFF0000000000000"),
-            ("min", "0x7FF0000000000000"),
-        ):
-            spec = Reduce2DSpec(n_per_block=4096, op=op, dtype="f64", vec=2)
-            ll = lower_kernel_to_llvm(build_reduce2d(spec))
-            self.assertIn(f"double {inf_bits}", ll)
-
     def test_f64_rejects_wide_vec(self):
         from rocke.instances.common.reduce import Reduce2DSpec, is_valid_spec
 

From bb4ceff753ed3f92aaed2f78256c0301e044a483 Mon Sep 17 00:00:00 2001
From: Minseob Shin <minseob.shin11@gmail.com>
Date: Sun, 11 Oct 2026 23:11:04 +0000
Subject: [PATCH 5/7] keep f64 out of io helper (io_ir_type) and resolve
 elementwise/reduce I/O type via dtype_to_ir_type

---
 .../rocke/platform/cpp/helpers/io.cpp             | 11 ++---------
 .../platform/cpp/instances/common/elementwise.cpp |  3 +--
 .../platform/cpp/instances/common/reduce.cpp      |  4 ++--
 .../rocke/platform/python/rocke/helpers/io.py     | 10 ++++------
 .../python/rocke/instances/common/elementwise.py  | 13 ++++++++++---
 .../python/rocke/instances/common/reduce.py       | 15 ++++++++++++---
 6 files changed, 31 insertions(+), 25 deletions(-)

diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/helpers/io.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/helpers/io.cpp
index d6bcf65e178c..e90e86ded9e2 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/helpers/io.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/helpers/io.cpp
@@ -11,16 +11,14 @@
  *     def io_ir_type(dtype: str) -> Type:
  *         if dtype in ("f16", "fp16"): return F16
  *         if dtype == "bf16":          return BF16
- *         if dtype == "f64":           return F64
  *         raise ValueError(
- *             f"unsupported I/O dtype {dtype!r}; expected f16/fp16/bf16/f64")
+ *             f"unsupported I/O dtype {dtype!r}; expected f16/fp16/bf16")
  *
  * Mapping invariants (must stay byte-identical to the Python so downstream IR is
  * identical):
  *   "f16"  -> rocke_f16()    (Python F16)
  *   "fp16" -> rocke_f16()    (alias; Python F16)
  *   "bf16" -> rocke_bf16()   (Python BF16)
- *   "f64"  -> rocke_f64()    (Python F64)
  *   else   -> NULL         (Python ValueError)
  */
 
@@ -54,11 +52,6 @@ const rocke_type_t* rocke_io_ir_type(const char* dtype)
     {
         return rocke_bf16();
     }
-    /* `dtype == "f64"` -> F64 */
-    if(strcmp(dtype, "f64") == 0)
-    {
-        return rocke_f64();
-    }
     /* Python: raise ValueError. No builder here, so signal via NULL. */
     return NULL;
 }
@@ -82,7 +75,7 @@ const rocke_type_t* rocke_b_io_ir_type(rocke_ir_builder_t* b, const char* dtype)
         return (const rocke_type_t*)rocke_i_set_err(
             b,
             ROCKE_ERR_VALUE,
-            "unsupported I/O dtype %s%s%s; expected f16/fp16/bf16/f64",
+            "unsupported I/O dtype %s%s%s; expected f16/fp16/bf16",
             dtype ? "'" : "",
             dtype ? dtype : "None",
             dtype ? "'" : "");
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/elementwise.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/elementwise.cpp
index fe1dbec08f3e..8d8c3aaa6030 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/elementwise.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/elementwise.cpp
@@ -28,7 +28,6 @@
 #include "rocke/error_boundary.hpp" /* ckc::guard_builder boundary shim */
 #include "rocke/helper_rocke.helpers.activations.h"
 #include "rocke/helper_rocke.helpers.distribution.h"
-#include "rocke/helper_rocke.helpers.io.h"
 #include "rocke/helper_rocke.helpers.spec.h"
 #include "rocke/helper_rocke.helpers.tensor_view.h"
 #include "rocke/ir_internal.h" /* rocke_i_set_err */
@@ -505,7 +504,7 @@ rocke_kernel_def_t* rocke_build_elementwise(rocke_ir_builder_t* b,
         }
     }
 
-    io_ty = rocke_b_io_ir_type(b, spec->dtype);
+    io_ty = rocke_dtype_to_ir_type(spec->dtype);
     if(io_ty == NULL)
     {
         return NULL;
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/reduce.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/reduce.cpp
index 1dc317c368c7..6b7e4f66641b 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/reduce.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/reduce.cpp
@@ -377,8 +377,8 @@ rocke_kernel_def_t*
             return NULL;
         }
 
-        /* io_ty = io_ir_type(spec.dtype) */
-        io_ty = rocke_b_io_ir_type(b, spec->dtype);
+        /* io_ty = dtype_to_ir_type(spec.dtype) */
+        io_ty = rocke_dtype_to_ir_type(spec->dtype);
         if(io_ty == NULL)
         {
             return NULL;
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/io.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/io.py
index 3ace81fa7e94..379484c5ab18 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/io.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/io.py
@@ -22,7 +22,7 @@
 
 from typing import Literal
 
-from ..core.ir import BF16, F16, F64, IRBuilder, TemporalHint, Type, Value
+from ..core.ir import BF16, F16, IRBuilder, TemporalHint, Type, Value
 
 
 __all__ = [
@@ -51,17 +51,15 @@
 def io_ir_type(dtype: str) -> Type:
     """Map a dtype string to the canonical IR type object.
 
-    Accepts ``"f16"``, ``"fp16"`` (alias), ``"bf16"``, or ``"f64"``.
-    Raises :class:`ValueError` for any other value -- f8 and i8 paths go
+    Accepts ``"f16"``, ``"fp16"`` (alias), or ``"bf16"``. Raises
+    :class:`ValueError` for any other value -- f8 and i8 paths go
     through their own helpers because their compute dtype isn't f32.
     """
     if dtype in ("f16", "fp16"):
         return F16
     if dtype == "bf16":
         return BF16
-    if dtype == "f64":
-        return F64
-    raise ValueError(f"unsupported I/O dtype {dtype!r}; expected f16/fp16/bf16/f64")
+    raise ValueError(f"unsupported I/O dtype {dtype!r}; expected f16/fp16/bf16")
 
 
 def load_scalar(b: IRBuilder, ptr: Value, idx: Value, *, dtype: str) -> Value:
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/common/elementwise.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/common/elementwise.py
index 5bdbee520d25..9e72033bb46a 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/common/elementwise.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/common/elementwise.py
@@ -52,14 +52,21 @@
 from dataclasses import dataclass
 from typing import Literal, Tuple
 
-from ...core.ir import I32, IRBuilder, KernelDef, PtrType, Type, Value
+from ...core.ir import (
+    I32,
+    IRBuilder,
+    KernelDef,
+    PtrType,
+    Type,
+    Value,
+    dtype_to_ir_type,
+)
 from ...helpers.activations import _sigmoid_via_exp2, _tanh_via_exp2
 from ...helpers.distribution import (
     TileDistributionEncoding,
     make_static_distributed_tensor,
     make_static_tile_distribution,
 )
-from ...helpers.io import io_ir_type
 from ...helpers.spec import SignatureBuilder, kernel_name_join
 from ...helpers.tensor_view import make_global_view, make_tile_window
 
@@ -272,7 +279,7 @@ def build_elementwise(spec: ElementwiseSpec) -> KernelDef:
     if not ok:
         raise ValueError(f"invalid elementwise spec: {why}")
 
-    io_ty = io_ir_type(spec.dtype)
+    io_ty = dtype_to_ir_type(spec.dtype)
 
     b = IRBuilder(spec.kernel_name())
     b.kernel.attrs["max_workgroup_size"] = spec.block_size
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/common/reduce.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/common/reduce.py
index 8facfb4350e2..44c08c7e2230 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/common/reduce.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/common/reduce.py
@@ -40,7 +40,16 @@
 from dataclasses import dataclass
 from typing import Literal, Tuple
 
-from ...core.ir import F32, F64, I32, IRBuilder, KernelDef, PtrType, Value
+from ...core.ir import (
+    F32,
+    F64,
+    I32,
+    IRBuilder,
+    KernelDef,
+    PtrType,
+    Value,
+    dtype_to_ir_type,
+)
 from ...helpers.distribution import (
     TileDistributionEncoding,
     block_tile_reduce_sync,
@@ -48,7 +57,7 @@
     make_static_distributed_tensor,
     make_static_tile_distribution,
 )
-from ...helpers.io import io_ir_type, store_scalar_from_f32
+from ...helpers.io import store_scalar_from_f32
 from ...helpers.reduction import (
     block_lds_reduce,
     block_lds_reduce_with_wave_prologue,
@@ -193,7 +202,7 @@ def build_reduce2d(spec: Reduce2DSpec) -> KernelDef:
     if not ok:
         raise ValueError(f"invalid reduce2d spec: {why}")
 
-    io_ty = io_ir_type(spec.dtype)
+    io_ty = dtype_to_ir_type(spec.dtype)
     BS, VEC, N = spec.block_size, spec.vec, spec.n_per_block
 
     b = IRBuilder(spec.kernel_name())

From 6b8fefdde2e002592db1eb9199198c15e31d0415 Mon Sep 17 00:00:00 2001
From: Minseob Shin <minseob.shin11@gmail.com>
Date: Tue, 13 Oct 2026 00:59:17 +0000
Subject: [PATCH 6/7] Regenerate intrinsic arch-domain artifacts for
 maxnum.f64/minnum.f64

---
 .../data/intrinsic_arch_domain.llvm20.json    | 66 ++++++++++++++++++-
 .../data/intrinsic_arch_domain.llvm22.json    | 64 +++++++++++++++++-
 .../data/intrinsic_arch_domain.llvm23.json    | 62 +++++++++++++++++
 3 files changed, 190 insertions(+), 2 deletions(-)

diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm20.json b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm20.json
index 59e66b458282..15e68c02ef53 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm20.json
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm20.json
@@ -50,6 +50,7 @@
     "maxnum.bf16": "llvm.maxnum.bf16",
     "maxnum.f16": "llvm.maxnum.f16",
     "maxnum.f32": "llvm.maxnum.f32",
+    "maxnum.f64": "llvm.maxnum.f64",
     "mbcnt.hi": "llvm.amdgcn.mbcnt.hi",
     "mbcnt.lo": "llvm.amdgcn.mbcnt.lo",
     "mfma.f32.16x16x128.fp8.hero": "llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32",
@@ -74,6 +75,7 @@
     "minnum.bf16": "llvm.minnum.bf16",
     "minnum.f16": "llvm.minnum.f16",
     "minnum.f32": "llvm.minnum.f32",
+    "minnum.f64": "llvm.minnum.f64",
     "mov.dpp8.f32": "llvm.amdgcn.mov.dpp8.f32",
     "mov.dpp8.i32": "llvm.amdgcn.mov.dpp8.i32",
     "raw.ptr.buffer.load.async.lds": "llvm.amdgcn.raw.ptr.buffer.load.async.lds",
@@ -2322,6 +2324,37 @@
         "verified_on": "llvm20"
       }
     },
+    "maxnum.f64": {
+      "gfx11-generic": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      },
+      "gfx1151": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      },
+      "gfx1201": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      },
+      "gfx1250": {
+        "evidence": "invalid target ID",
+        "status": "target_unsupported",
+        "verified_on": "llvm20"
+      },
+      "gfx90a": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      },
+      "gfx942": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      },
+      "gfx950": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      }
+    },
     "mbcnt.hi": {
       "gfx11-generic": {
         "status": "ok",
@@ -3217,6 +3250,37 @@
         "verified_on": "llvm20"
       }
     },
+    "minnum.f64": {
+      "gfx11-generic": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      },
+      "gfx1151": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      },
+      "gfx1201": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      },
+      "gfx1250": {
+        "evidence": "invalid target ID",
+        "status": "target_unsupported",
+        "verified_on": "llvm20"
+      },
+      "gfx90a": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      },
+      "gfx942": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      },
+      "gfx950": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      }
+    },
     "mov.dpp8.f32": {
       "gfx11-generic": {
         "status": "ok",
@@ -5865,7 +5929,7 @@
       "gfx942",
       "gfx950"
     ],
-    "clang": "AMD clang version 20.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-7.1.1 25444 27682a16360e33e37c4f3cc6adf9a620733f8fe1)",
+    "clang": "AMD clang version 20.0.0git (https://github.com/ROCm/llvm-project.git 27682a16360e33e37c4f3cc6adf9a620733f8fe1)",
     "flavor": "llvm20",
     "generator": 4,
     "probe_cflags": [
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm22.json b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm22.json
index f2d74722cd2e..f56744efb1ca 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm22.json
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm22.json
@@ -60,6 +60,7 @@
     "maxnum.bf16": "llvm.maxnum.bf16",
     "maxnum.f16": "llvm.maxnum.f16",
     "maxnum.f32": "llvm.maxnum.f32",
+    "maxnum.f64": "llvm.maxnum.f64",
     "mbcnt.hi": "llvm.amdgcn.mbcnt.hi",
     "mbcnt.lo": "llvm.amdgcn.mbcnt.lo",
     "mfma.f32.16x16x128.fp8.hero": "llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32",
@@ -83,6 +84,7 @@
     "minnum.bf16": "llvm.minnum.bf16",
     "minnum.f16": "llvm.minnum.f16",
     "minnum.f32": "llvm.minnum.f32",
+    "minnum.f64": "llvm.minnum.f64",
     "mov.dpp8.f32": "llvm.amdgcn.mov.dpp8.f32",
     "mov.dpp8.i32": "llvm.amdgcn.mov.dpp8.i32",
     "raw.ptr.buffer.load.i16": "llvm.amdgcn.raw.ptr.buffer.load.i16",
@@ -2271,6 +2273,36 @@
         "verified_on": "llvm22"
       }
     },
+    "maxnum.f64": {
+      "gfx11-generic": {
+        "status": "ok",
+        "verified_on": "llvm22"
+      },
+      "gfx1151": {
+        "status": "ok",
+        "verified_on": "llvm22"
+      },
+      "gfx1201": {
+        "status": "ok",
+        "verified_on": "llvm22"
+      },
+      "gfx1250": {
+        "status": "ok",
+        "verified_on": "llvm22"
+      },
+      "gfx90a": {
+        "status": "ok",
+        "verified_on": "llvm22"
+      },
+      "gfx942": {
+        "status": "ok",
+        "verified_on": "llvm22"
+      },
+      "gfx950": {
+        "status": "ok",
+        "verified_on": "llvm22"
+      }
+    },
     "mbcnt.hi": {
       "gfx11-generic": {
         "status": "ok",
@@ -3162,6 +3194,36 @@
         "verified_on": "llvm22"
       }
     },
+    "minnum.f64": {
+      "gfx11-generic": {
+        "status": "ok",
+        "verified_on": "llvm22"
+      },
+      "gfx1151": {
+        "status": "ok",
+        "verified_on": "llvm22"
+      },
+      "gfx1201": {
+        "status": "ok",
+        "verified_on": "llvm22"
+      },
+      "gfx1250": {
+        "status": "ok",
+        "verified_on": "llvm22"
+      },
+      "gfx90a": {
+        "status": "ok",
+        "verified_on": "llvm22"
+      },
+      "gfx942": {
+        "status": "ok",
+        "verified_on": "llvm22"
+      },
+      "gfx950": {
+        "status": "ok",
+        "verified_on": "llvm22"
+      }
+    },
     "mov.dpp8.f32": {
       "gfx11-generic": {
         "status": "ok",
@@ -5727,7 +5789,7 @@
       "gfx942",
       "gfx950"
     ],
-    "clang": "AMD clang version 22.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-7.2.0 26014 7b800a19466229b8479a78de19143dc33c3ab9b5)",
+    "clang": "AMD clang version 22.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-7.2.3 26084 f58b06dce1f9c15707c5f808fd002e18c2accf7e)",
     "flavor": "llvm22",
     "generator": 4,
     "probe_cflags": [
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm23.json b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm23.json
index a8fbafec0ea9..3277414f7d66 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm23.json
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm23.json
@@ -65,6 +65,7 @@
     "maxnum.bf16": "llvm.maxnum.bf16",
     "maxnum.f16": "llvm.maxnum.f16",
     "maxnum.f32": "llvm.maxnum.f32",
+    "maxnum.f64": "llvm.maxnum.f64",
     "mbcnt.hi": "llvm.amdgcn.mbcnt.hi",
     "mbcnt.lo": "llvm.amdgcn.mbcnt.lo",
     "mfma.f32.16x16x128.fp8.hero": "llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32",
@@ -89,6 +90,7 @@
     "minnum.bf16": "llvm.minnum.bf16",
     "minnum.f16": "llvm.minnum.f16",
     "minnum.f32": "llvm.minnum.f32",
+    "minnum.f64": "llvm.minnum.f64",
     "mov.dpp8.f32": "llvm.amdgcn.mov.dpp8.f32",
     "mov.dpp8.i32": "llvm.amdgcn.mov.dpp8.i32",
     "raw.ptr.buffer.load.async.lds": "llvm.amdgcn.raw.ptr.buffer.load.async.lds",
@@ -2254,6 +2256,36 @@
         "verified_on": "llvm23"
       }
     },
+    "maxnum.f64": {
+      "gfx11-generic": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      },
+      "gfx1151": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      },
+      "gfx1201": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      },
+      "gfx1250": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      },
+      "gfx90a": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      },
+      "gfx942": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      },
+      "gfx950": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      }
+    },
     "mbcnt.hi": {
       "gfx11-generic": {
         "status": "ok",
@@ -3144,6 +3176,36 @@
         "verified_on": "llvm23"
       }
     },
+    "minnum.f64": {
+      "gfx11-generic": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      },
+      "gfx1151": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      },
+      "gfx1201": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      },
+      "gfx1250": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      },
+      "gfx90a": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      },
+      "gfx942": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      },
+      "gfx950": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      }
+    },
     "mov.dpp8.f32": {
       "gfx11-generic": {
         "status": "ok",

From c1d5eefd3c30fe3f455ddae766ba1e52f1f22e71 Mon Sep 17 00:00:00 2001
From: Minseob Shin <minseob.shin11@gmail.com>
Date: Tue, 13 Oct 2026 21:46:59 +0000
Subject: [PATCH 7/7] Reject unmappable dtypes in fusion lowering and use
 fmax/fmin numeric references

---
 .../platform/python/rocke/helpers/fusion_lowering.py   |  6 ++++++
 .../platform/tests/instances/differential/numeric.py   | 10 +++++-----
 2 files changed, 11 insertions(+), 5 deletions(-)

diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/fusion_lowering.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/fusion_lowering.py
index 5c9e8dd41ccd..ec8eaf79e97a 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/fusion_lowering.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/fusion_lowering.py
@@ -571,6 +571,9 @@ def can_lower(self, graph: FusionGraph, region: FusionRegion) -> LegalResult:
         op = graph.ops[region.op_names[0]]
         if op.kind not in (self._UNARY | self._BINARY):
             return LegalResult.failure(f"{self.name}: unsupported op {op.kind!r}")
+        dtype = _dtype_str_for_tensor(graph, op.inputs[0])
+        if dtype not in ("fp16", "bf16"):
+            return LegalResult.failure(f"{self.name}: unsupported dtype {dtype!r}")
         return LegalResult.success()
 
     def candidates(
@@ -672,6 +675,9 @@ def can_lower(self, graph: FusionGraph, region: FusionRegion) -> LegalResult:
         in_t = graph.tensors[op.inputs[0]]
         if len(in_t.shape) != 2:
             return LegalResult.failure(f"{self.name}: input must be 2D")
+        dtype = _dtype_str_for_tensor(graph, op.inputs[0])
+        if dtype not in ("fp16", "bf16"):
+            return LegalResult.failure(f"{self.name}: unsupported dtype {dtype!r}")
         return LegalResult.success()
 
     def candidates(
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/numeric.py b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/numeric.py
index 4bb19ca3f1e1..53db4eb9b75e 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/numeric.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/numeric.py
@@ -577,11 +577,11 @@ def _torch_elementwise_ref(torch, op: str, A, B):
     if op == "mul":
         return a * B.to(rd)
     if op == "max":
-        return torch.maximum(a, B.to(rd))
+        return torch.fmax(a, B.to(rd))
     if op == "min":
-        return torch.minimum(a, B.to(rd))
+        return torch.fmin(a, B.to(rd))
     if op == "relu":
-        return torch.relu(a)
+        return torch.fmax(a, torch.zeros_like(a))
     if op == "silu":
         return a * torch.sigmoid(a)
     if op == "neg":
@@ -1121,9 +1121,9 @@ def run_row_config(cfg: RowCfg, arch: str = "gfx950") -> NumericResult:
         if cfg.op == "sum":
             ref = x_ref.sum(dim=-1)
         elif cfg.op == "max":
-            ref = x_ref.max(dim=-1).values
+            ref = torch.fmax(x_ref, torch.full_like(x_ref, float("-inf"))).amax(dim=-1)
         elif cfg.op == "min":
-            ref = x_ref.min(dim=-1).values
+            ref = torch.fmin(x_ref, torch.full_like(x_ref, float("inf"))).amin(dim=-1)
         elif cfg.op == "mean":
             ref = x_ref.mean(dim=-1)
         else:  # pragma: no cover