replace UNUSED with std::ignore

include/ck/config.hpp

@@ -135,9 +135,6 @@
 // https://github.com/ROCmSoftwarePlatform/composable_kernel/issues/135
 #define CK_WORKAROUND_GITHUB_135 1
 
-//workaround for building xdl kernels on all architectures
-#define UNUSED(x) (void)(x)
-
 namespace ck {
 
 enum struct InMemoryDataOperationEnum

include/ck/tensor_operation/gpu/device/device_batched_gemm_reduce_xdl_cshuffle.hpp

@@ -90,23 +90,23 @@ __global__ void
                                                    d_grid_desc_mblock_mperblock,
                                                    block_2_ctile_map);
 #else
-	UNUSED(p_a_grid);
-	UNUSED(p_b_grid);
-	UNUSED(p_c_grid);
-	UNUSED(p_d0_grid);
-	UNUSED(p_d1_grid);
-	UNUSED(batch_count);
-	UNUSED(a_element_op);
-	UNUSED(b_element_op);
-	UNUSED(c_element_op);
-	UNUSED(d0_reduce_op);
-	UNUSED(d1_reduce_op);
-	UNUSED(a_grid_desc_ak0_m_ak1);
-	UNUSED(b_grid_desc_bk0_n_bk1);
-	UNUSED(c_grid_desc_mblock_mperblock_nblock_nperblock);
-	UNUSED(d_grid_desc_mblock_mperblock);
-	UNUSED(compute_base_ptr_of_batch_);
-	UNUSED(block_2_ctile_map);
+	std::ignore = p_a_grid;
+	std::ignore = p_b_grid;
+	std::ignore = p_c_grid;
+	std::ignore = p_d0_grid;
+	std::ignore = p_d1_grid;
+	std::ignore = batch_count;
+	std::ignore = a_element_op;
+	std::ignore = b_element_op;
+	std::ignore = c_element_op;
+	std::ignore = d0_reduce_op;
+	std::ignore = d1_reduce_op;
+	std::ignore = a_grid_desc_ak0_m_ak1;
+	std::ignore = b_grid_desc_bk0_n_bk1;
+	std::ignore = c_grid_desc_mblock_mperblock_nblock_nperblock;
+	std::ignore = d_grid_desc_mblock_mperblock;
+	std::ignore = compute_base_ptr_of_batch_;
+	std::ignore = block_2_ctile_map;
 #endif //end of if defined (defined(__gfx908__) || defined(__gfx90a__))												   
 }
 

include/ck/tensor_operation/gpu/device/device_batched_gemm_xdl.hpp

@@ -72,18 +72,18 @@ __global__ void
                                                   c_element_op,
                                                   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);
+	std::ignore = p_a_grid;
+	std::ignore = p_b_grid;
+	std::ignore = p_c_grid;
+	std::ignore = batch_count;
+	std::ignore = a_grid_desc_k0_m_k1;
+	std::ignore = b_grid_desc_k0_n_k1;
+	std::ignore = c_grid_desc_m0_n0_m1_n1_m2_m3_m4_n2;
+	std::ignore = a_element_op;
+	std::ignore = b_element_op;
+	std::ignore = c_element_op;
+	std::ignore = compute_base_ptr_of_batch_;
+	std::ignore = block_2_ctile_map;
 #endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
 }
 

include/ck/tensor_operation/gpu/device/device_conv3d_fwd_xdl_ndhwc_kzyxc_ndhwk.hpp

@@ -76,20 +76,20 @@ __global__ void
                                                   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);
+	std::ignore = p_a_grid;
+	std::ignore = p_b_grid;
+	std::ignore = p_c_grid;
+	std::ignore = num_batches;
+	std::ignore = a_batch_stride;
+	std::ignore = b_batch_stride;
+	std::ignore = c_batch_stride;
+	std::ignore = a_grid_desc_k0_m_k1;
+	std::ignore = b_grid_desc_k0_n_k1;
+	std::ignore = c_grid_desc_m0_n0_m1_n1_m2_m3_m4_n2;
+	std::ignore = a_element_op;
+	std::ignore = b_element_op;
+	std::ignore = c_element_op;
+	std::ignore = block_2_ctile_map;
 #endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
 }
 

include/ck/tensor_operation/gpu/grid/gridwise_gemm_reduce_xdl_cshuffle_v1.hpp

@@ -68,22 +68,21 @@ __global__ void
                                                    d_grid_desc_mblock_mperblock,
                                                    block_2_ctile_map);
 #else
-	UNUSED(p_a_grid);
-	UNUSED(p_b_grid);
-	UNUSED(p_c_grid);
-	UNUSED(p_d0_grid);
-	UNUSED(p_d1_grid);
-	UNUSED(a_element_op);
-	UNUSED(b_element_op);
-	UNUSED(c_element_op);
-	UNUSED(d0_reduce_op);
-	UNUSED(d1_reduce_op);
-	UNUSED(a_grid_desc_ak0_m_ak1);
-	UNUSED(b_grid_desc_bk0_n_bk1);
-	UNUSED(c_grid_desc_mblock_mperblock_nblock_nperblock);
-	UNUSED(d_grid_desc_mblock_mperblock);	
-	UNUSED(block_2_ctile_map);	
-	
+	std::ignore = p_a_grid;
+	std::ignore = p_b_grid;
+	std::ignore = p_c_grid;
+	std::ignore = p_d0_grid;
+	std::ignore = p_d1_grid;
+	std::ignore = a_element_op;
+	std::ignore = b_element_op;
+	std::ignore = c_element_op;
+	std::ignore = d0_reduce_op;
+	std::ignore = d1_reduce_op;
+	std::ignore = a_grid_desc_ak0_m_ak1;
+	std::ignore = b_grid_desc_bk0_n_bk1;
+	std::ignore = c_grid_desc_mblock_mperblock_nblock_nperblock;
+	std::ignore = d_grid_desc_mblock_mperblock;	
+	std::ignore = block_2_ctile_map;	
 #endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
 }
 

include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v1.hpp

@@ -53,16 +53,16 @@ __global__ void
                                                    c_grid_desc_mblock_mperblock_nblock_nperblock,
                                                    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);	
+	std::ignore = p_a_grid;
+	std::ignore = p_b_grid;
+	std::ignore = p_c_grid;
+	std::ignore = a_element_op;
+	std::ignore = b_element_op;
+	std::ignore = c_element_op;
+	std::ignore = a_grid_desc_ak0_m_ak1;
+	std::ignore = b_grid_desc_bk0_n_bk1;
+	std::ignore = c_grid_desc_mblock_mperblock_nblock_nperblock;
+	std::ignore = block_2_ctile_map;	
 #endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
 }
 

include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r3.hpp

@@ -54,16 +54,16 @@ __global__ void
                                                    c_element_op,
                                                    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);	
+	std::ignore = p_a_grid;
+	std::ignore = p_b_grid;
+	std::ignore = p_c_grid;
+	std::ignore = a_grid_desc_k0_m_k1;
+	std::ignore = b_grid_desc_k0_n_k1;
+	std::ignore = c_grid_desc_m0_n0_m1_n1_m2_m3_m4_n2;
+	std::ignore = a_element_op;
+	std::ignore = b_element_op;
+	std::ignore = c_element_op;
+	std::ignore = block_2_ctile_map;	
 #endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
 }
 
@@ -141,11 +141,11 @@ __global__ void
         block_id_grp);
 #endif
 #else
-	UNUSED(gemm_desc_);
-	UNUSED(group_count);
-	UNUSED(a_element_op);
-	UNUSED(b_element_op);
-	UNUSED(c_element_op);
+	std::ignore = gemm_desc_;
+	std::ignore = group_count;
+	std::ignore = a_element_op;
+	std::ignore = b_element_op;
+	std::ignore = c_element_op;
 #endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
 }
 

include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r4.hpp

@@ -55,16 +55,16 @@ __global__ void
                                                   c_element_op,
                                                   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);	
+	std::ignore = p_a_grid;
+	std::ignore = p_b_grid;
+	std::ignore = p_c_grid;
+	std::ignore = a_b_k0_m_k1_grid_desc;
+	std::ignore = b_b_k0_n_k1_grid_desc;
+	std::ignore = c_m0_n0_m1_n1_m2_m3_m4_n2_grid_desc;
+	std::ignore = a_element_op;
+	std::ignore = b_element_op;
+	std::ignore = c_element_op;
+	std::ignore = c_block_cluster_adaptor);	
 #endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
 }
 

include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r4r2.hpp

@@ -57,16 +57,16 @@ __global__ void
                                                   c_element_op,
                                                   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);
+	std::ignore = p_a_grid;
+	std::ignore = p_b_grid;
+	std::ignore = p_c_grid;
+	std::ignore = a_b_k0_m_k1_grid_desc;
+	std::ignore = b_b_k0_n_k1_grid_desc;
+	std::ignore = c_grid_desc_mblock_mperblock_nblock_nperblock;
+	std::ignore = a_element_op;
+	std::ignore = b_element_op;
+	std::ignore = c_element_op;
+	std::ignore = c_block_cluster_adaptor;
 #endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
 }
 

include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v3r1.hpp

@@ -58,16 +58,16 @@ __global__ void
         c_element_op,
         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);
+	std::ignore = p_a_grid;
+	std::ignore = p_b_grid;
+	std::ignore = p_c_grid;
+	std::ignore = a_grid_desc_ak0_m_ak1;
+	std::ignore = b_grid_desc_bk0_n_bk1;
+	std::ignore = c_grid_desc_mblock_mxdlperwave_mwavemperxdl_nblock_nxdlperwave_nwavenperxdl;
+	std::ignore = a_element_op;
+	std::ignore = b_element_op;
+	std::ignore = c_element_op;
+	std::ignore = block_2_ctile_map;
 #endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
 }
 

include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v3r2.hpp

@@ -63,18 +63,18 @@ __global__ void
         c_element_op,
         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);
+	std::ignore = p_a_grid;
+	std::ignore = p_b_grid;
+	std::ignore = p_c_grid;
+	std::ignore = p_c0_grid;
+	std::ignore = a_grid_desc_k0_m_k1;
+	std::ignore = b_grid_desc_k0_n_k1;
+	std::ignore = c_grid_desc_mblock_mxdlperwave_mwavemperxdl_nblock_nxdlperwave_nwavenperxdl;
+	std::ignore = c0_grid_desc_mblock_mxdlperwave_mwavemperxdl_nblock_nxdlperwave_nwavenperxdl;
+	std::ignore = a_element_op;
+	std::ignore = b_element_op;
+	std::ignore = c_element_op;
+	std::ignore = block_2_ctile_map;
 #endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
 }
 

include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v3r3.hpp

@@ -69,20 +69,20 @@ __global__ void
         c_element_op,
         block_2_ctile_map);
 #else
-	UNUSED(p_a_grid);
-	UNUSED(p_b_grid);
-	UNUSED(p_c_grid);
-	UNUSED(p_c0_grid);
-	UNUSED(p_c1_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(c1_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);
+	std::ignore = p_a_grid;
+	std::ignore = p_b_grid;
+	std::ignore = p_c_grid;
+	std::ignore = p_c0_grid;
+	std::ignore = p_c1_grid;
+	std::ignore = a_grid_desc_k0_m_k1;
+	std::ignore = b_grid_desc_k0_n_k1;
+	std::ignore = c_grid_desc_mblock_mxdlperwave_mwavemperxdl_nblock_nxdlperwave_nwavenperxdl;
+	std::ignore = c0_grid_desc_mblock_mxdlperwave_mwavemperxdl_nblock_nxdlperwave_nwavenperxdl;
+	std::ignore = c1_grid_desc_mblock_mxdlperwave_mwavemperxdl_nblock_nxdlperwave_nwavenperxdl;
+	std::ignore = a_element_op;
+	std::ignore = b_element_op;
+	std::ignore = c_element_op;
+	std::ignore = block_2_ctile_map;
 #endif //end of if (defined(__gfx908__) || defined(__gfx90a__))
 }