Skip to content
GitLab
Menu
Projects
Groups
Snippets
Loading...
Help
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in / Register
Toggle navigation
Menu
Open sidebar
gaoqiong
composable_kernel
Commits
b5e62fde
Commit
b5e62fde
authored
Apr 11, 2022
by
illsilin
Browse files
test the warning suppression
parent
1a685834
Changes
8
Hide whitespace changes
Inline
Side-by-side
Showing
8 changed files
with
104 additions
and
1 deletion
+104
-1
include/ck/tensor_operation/gpu/device/device_batched_gemm_xdl.hpp
...k/tensor_operation/gpu/device/device_batched_gemm_xdl.hpp
+14
-1
include/ck/tensor_operation/gpu/device/device_conv3d_fwd_xdl_ndhwc_kzyxc_ndhwk.hpp
...on/gpu/device/device_conv3d_fwd_xdl_ndhwc_kzyxc_ndhwk.hpp
+16
-0
include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v1.hpp
...nsor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v1.hpp
+11
-0
include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r3.hpp
...k/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r3.hpp
+17
-0
include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r4.hpp
...k/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r4.hpp
+11
-0
include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r4r2.hpp
...tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r4r2.hpp
+11
-0
include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v3r1.hpp
...k/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v3r1.hpp
+11
-0
include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v3r2.hpp
...k/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v3r2.hpp
+13
-0
No files found.
include/ck/tensor_operation/gpu/device/device_batched_gemm_xdl.hpp
View file @
b5e62fde
...
@@ -71,6 +71,20 @@ __global__ void
...
@@ -71,6 +71,20 @@ __global__ void
b_element_op
,
b_element_op
,
c_element_op
,
c_element_op
,
block_2_ctile_map
);
block_2_ctile_map
);
#else
UNUSED
(
p_a_grid
);
UNUSED
(
p_b_grid
);
UNUSED
(
p_c_grid
);
UNUSED
(
batch_count
);
UNUSED
(
a_grid_desc_k0_m_k1
);
UNUSED
(
b_grid_desc_k0_n_k1
);
UNUSED
(
c_grid_desc_m0_n0_m1_n1_m2_m3_m4_n2
);
UNUSED
(
a_element_op
);
UNUSED
(
b_element_op
);
UNUSED
(
c_element_op
);
UNUSED
(
compute_base_ptr_of_batch_
);
UNUSED
(
block_2_ctile_map
);
#endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
}
}
template
<
typename
ADataType
,
template
<
typename
ADataType
,
...
@@ -610,7 +624,6 @@ struct DeviceBatchedGemmXdl
...
@@ -610,7 +624,6 @@ struct DeviceBatchedGemmXdl
return
str
.
str
();
return
str
.
str
();
}
}
#endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
};
};
}
// namespace device
}
// namespace device
...
...
include/ck/tensor_operation/gpu/device/device_conv3d_fwd_xdl_ndhwc_kzyxc_ndhwk.hpp
View file @
b5e62fde
...
@@ -74,6 +74,22 @@ __global__ void
...
@@ -74,6 +74,22 @@ __global__ void
b_element_op
,
b_element_op
,
c_element_op
,
c_element_op
,
block_2_ctile_map
);
block_2_ctile_map
);
#else
UNUSED
(
p_a_grid
);
UNUSED
(
p_b_grid
);
UNUSED
(
p_c_grid
);
UNUSED
(
num_batches
);
UNUSED
(
a_batch_stride
);
UNUSED
(
b_batch_stride
);
UNUSED
(
c_batch_stride
);
UNUSED
(
a_grid_desc_k0_m_k1
);
UNUSED
(
b_grid_desc_k0_n_k1
);
UNUSED
(
c_grid_desc_m0_n0_m1_n1_m2_m3_m4_n2
);
UNUSED
(
a_element_op
);
UNUSED
(
b_element_op
);
UNUSED
(
c_element_op
);
UNUSED
(
block_2_ctile_map
);
#endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
#endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
}
}
...
...
include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v1.hpp
View file @
b5e62fde
...
@@ -52,6 +52,17 @@ __global__ void
...
@@ -52,6 +52,17 @@ __global__ void
b_grid_desc_bk0_n_bk1
,
b_grid_desc_bk0_n_bk1
,
c_grid_desc_mblock_mperblock_nblock_nperblock
,
c_grid_desc_mblock_mperblock_nblock_nperblock
,
block_2_ctile_map
);
block_2_ctile_map
);
#else
UNUSED
(
p_a_grid
);
UNUSED
(
p_b_grid
);
UNUSED
(
p_c_grid
);
UNUSED
(
a_element_op
);
UNUSED
(
b_element_op
);
UNUSED
(
c_element_op
);
UNUSED
(
a_grid_desc_ak0_m_ak1
);
UNUSED
(
b_grid_desc_bk0_n_bk1
);
UNUSED
(
c_grid_desc_mblock_mperblock_nblock_nperblock
);
UNUSED
(
block_2_ctile_map
);
#endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
#endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
}
}
...
...
include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r3.hpp
View file @
b5e62fde
...
@@ -53,6 +53,17 @@ __global__ void
...
@@ -53,6 +53,17 @@ __global__ void
b_element_op
,
b_element_op
,
c_element_op
,
c_element_op
,
block_2_ctile_map
);
block_2_ctile_map
);
#else
UNUSED
(
p_a_grid
);
UNUSED
(
p_b_grid
);
UNUSED
(
p_c_grid
);
UNUSED
(
a_grid_desc_k0_m_k1
);
UNUSED
(
b_grid_desc_k0_n_k1
);
UNUSED
(
c_grid_desc_m0_n0_m1_n1_m2_m3_m4_n2
);
UNUSED
(
a_element_op
);
UNUSED
(
b_element_op
);
UNUSED
(
c_element_op
);
UNUSED
(
block_2_ctile_map
);
#endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
#endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
}
}
...
@@ -129,6 +140,12 @@ __global__ void
...
@@ -129,6 +140,12 @@ __global__ void
gemm_desc_ptr
[
group_id
].
block_2_ctile_map_
,
gemm_desc_ptr
[
group_id
].
block_2_ctile_map_
,
block_id_grp
);
block_id_grp
);
#endif
#endif
#else
UNUSED
(
gemm_desc_
);
UNUSED
(
group_count
);
UNUSED
(
a_element_op
);
UNUSED
(
b_element_op
);
UNUSED
(
c_element_op
);
#endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
#endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
}
}
...
...
include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r4.hpp
View file @
b5e62fde
...
@@ -54,6 +54,17 @@ __global__ void
...
@@ -54,6 +54,17 @@ __global__ void
b_element_op
,
b_element_op
,
c_element_op
,
c_element_op
,
c_block_cluster_adaptor
);
c_block_cluster_adaptor
);
#else
UNUSED
(
p_a_grid
);
UNUSED
(
p_b_grid
);
UNUSED
(
p_c_grid
);
UNUSED
(
a_b_k0_m_k1_grid_desc
);
UNUSED
(
b_b_k0_n_k1_grid_desc
);
UNUSED
(
c_m0_n0_m1_n1_m2_m3_m4_n2_grid_desc
);
UNUSED
(
a_element_op
);
UNUSED
(
b_element_op
);
UNUSED
(
c_element_op
);
UNUSED
(
c_block_cluster_adaptor
);
#endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
#endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
}
}
...
...
include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r4r2.hpp
View file @
b5e62fde
...
@@ -56,6 +56,17 @@ __global__ void
...
@@ -56,6 +56,17 @@ __global__ void
b_element_op
,
b_element_op
,
c_element_op
,
c_element_op
,
c_block_cluster_adaptor
);
c_block_cluster_adaptor
);
#else
UNUSED
(
p_a_grid
);
UNUSED
(
p_b_grid
);
UNUSED
(
p_c_grid
);
UNUSED
(
a_b_k0_m_k1_grid_desc
);
UNUSED
(
b_b_k0_n_k1_grid_desc
);
UNUSED
(
c_grid_desc_mblock_mperblock_nblock_nperblock
);
UNUSED
(
a_element_op
);
UNUSED
(
b_element_op
);
UNUSED
(
c_element_op
);
UNUSED
(
c_block_cluster_adaptor
);
#endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
#endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
}
}
...
...
include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v3r1.hpp
View file @
b5e62fde
...
@@ -57,6 +57,17 @@ __global__ void
...
@@ -57,6 +57,17 @@ __global__ void
b_element_op
,
b_element_op
,
c_element_op
,
c_element_op
,
block_2_ctile_map
);
block_2_ctile_map
);
#else
UNUSED
(
p_a_grid
);
UNUSED
(
p_b_grid
);
UNUSED
(
p_c_grid
);
UNUSED
(
a_grid_desc_ak0_m_ak1
);
UNUSED
(
b_grid_desc_bk0_n_bk1
);
UNUSED
(
c_grid_desc_mblock_mxdlperwave_mwavemperxdl_nblock_nxdlperwave_nwavenperxdl
);
UNUSED
(
a_element_op
);
UNUSED
(
b_element_op
);
UNUSED
(
c_element_op
);
UNUSED
(
block_2_ctile_map
);
#endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
#endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
}
}
...
...
include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v3r2.hpp
View file @
b5e62fde
...
@@ -62,6 +62,19 @@ __global__ void
...
@@ -62,6 +62,19 @@ __global__ void
b_element_op
,
b_element_op
,
c_element_op
,
c_element_op
,
block_2_ctile_map
);
block_2_ctile_map
);
#else
UNUSED
(
p_a_grid
);
UNUSED
(
p_b_grid
);
UNUSED
(
p_c_grid
);
UNUSED
(
p_c0_grid
);
UNUSED
(
a_grid_desc_k0_m_k1
);
UNUSED
(
b_grid_desc_k0_n_k1
);
UNUSED
(
c_grid_desc_mblock_mxdlperwave_mwavemperxdl_nblock_nxdlperwave_nwavenperxdl
);
UNUSED
(
c0_grid_desc_mblock_mxdlperwave_mwavemperxdl_nblock_nxdlperwave_nwavenperxdl
);
UNUSED
(
a_element_op
);
UNUSED
(
b_element_op
);
UNUSED
(
c_element_op
);
UNUSED
(
block_2_ctile_map
);
#endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
#endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
}
}
...
...
Write
Preview
Markdown
is supported
0%
Try again
or
attach a new file
.
Attach a file
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Cancel
Please
register
or
sign in
to comment