diff --git a/composable_kernel/include/utility/config_amd.hpp.in b/composable_kernel/include/utility/config_amd.hpp.in
index 071d48924c..437ed3ee8f 100644
--- a/composable_kernel/include/utility/config_amd.hpp.in
+++ b/composable_kernel/include/utility/config_amd.hpp.in
@@ -19,8 +19,6 @@ namespace ck {
 enum address_space_t
 {
     generic = 0,
-    vgpr    = 1,
-    lds     = 2,
     global  = 3
 };
 
diff --git a/composable_kernel/include/utility/config_nvidia.hpp.in b/composable_kernel/include/utility/config_nvidia.hpp.in
index ec5365a6ec..0844576ab8 100644
--- a/composable_kernel/include/utility/config_nvidia.hpp.in
+++ b/composable_kernel/include/utility/config_nvidia.hpp.in
@@ -21,8 +21,6 @@ namespace ck {
 enum address_space_t
 {
     generic = 0,
-    vgpr    = generic,
-    lds     = generic,
     global  = generic
 };
 
@@ -39,6 +37,17 @@ using index_t = int32_t;
 using float2_t = float2;
 using float4_t = float4;
 
+// data type conversion
+template <typename T>
+struct type_convert
+{
+    template <typename X>
+    __device__ T operator()(const X& x) const
+    {
+        return static_cast<T>(x);
+    }
+};
+
 template <class T>
 __device__ void fused_multiply_accumulate(T& d, const T& s0, const T& s1)
 {