diff --git a/include/ck_tile/core/arch/arch.hpp b/include/ck_tile/core/arch/arch.hpp
index 521d285739..a8237068d8 100644
--- a/include/ck_tile/core/arch/arch.hpp
+++ b/include/ck_tile/core/arch/arch.hpp
@@ -246,4 +246,13 @@ struct gfx12_t
 {
 };
 
+CK_TILE_DEVICE static constexpr auto get_device_ip()
+{
+#if defined(__gfx11__)
+    return gfx11_t{};
+#else // if defined(__gfx12__)
+    return gfx12_t{};
+#endif
+}
+
 } // namespace ck_tile
diff --git a/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma_impl.hpp b/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma_impl.hpp
index 36f757de63..623d7c5313 100644
--- a/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma_impl.hpp
+++ b/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma_impl.hpp
@@ -76,12 +76,9 @@ struct WarpGemmAttributeWmmaImpl
     }
 };
 
+using DeviceIp = remove_cvref_t<decltype(ck_tile::get_device_ip())>;
 using WarpGemmAttributeWmmaImpl_f32_16x16x16_f16_f16 = WarpGemmAttributeWmmaImpl<WmmaTraits<
-#if defined(__gfx11__)
-    gfx11_t,
-#else // if defined(__gfx12__)
-    gfx12_t,
-#endif
+    DeviceIp,
     fp16_t,
     fp16_t,
     float,
@@ -90,11 +87,7 @@ using WarpGemmAttributeWmmaImpl_f32_16x16x16_f16_f16 = WarpGemmAttributeWmmaImpl
     16>>;
 
 using WarpGemmAttributeWmmaImpl_f32_16x16x16_bf16_bf16 = WarpGemmAttributeWmmaImpl<WmmaTraits<
-#if defined(__gfx11__)
-    gfx11_t,
-#else // if defined(__gfx12__)
-    gfx12_t,
-#endif
+    DeviceIp,
     bf16_t,
     bf16_t,
     float,
@@ -103,11 +96,7 @@ using WarpGemmAttributeWmmaImpl_f32_16x16x16_bf16_bf16 = WarpGemmAttributeWmmaIm
     16>>;
 
 using WarpGemmAttributeWmmaImpl_i32_16x16x16_i8_i8 = WarpGemmAttributeWmmaImpl<WmmaTraits<
-#if defined(__gfx11__)
-    gfx11_t,
-#else // if defined(__gfx12__)
-    gfx12_t,
-#endif
+    DeviceIp,
     int8_t,
     int8_t,
     int32_t,