fixed

client_example/16_convnd_fwd/CMakeLists.txt

@@ -3,3 +3,6 @@ add_executable(client_conv3d_fwd_fp32 conv3d_fwd_fp32.cpp)
 
 target_link_libraries(client_conv3d_fwd_fp16 PRIVATE composable_kernel::device_operations)
 target_link_libraries(client_conv3d_fwd_fp32 PRIVATE composable_kernel::device_operations)
+
+add_executable(client_conv3d_fwd_fp16_comp_fp8 conv3d_fwd_fp16_comp_fp8.cpp)
+target_link_libraries(client_conv3d_fwd_fp16_comp_fp8 PRIVATE composable_kernel::device_operations)

client_example/16_convnd_fwd/common.hpp

@@ -95,7 +95,7 @@ template <ck::index_t NumDimSpatial,
           typename WeiLayout,
           typename OutLayout,
           ck::index_t NumNonSpatialDim = 3,
-          typename ComputeType = InDataType>
+          typename ComputeType         = InDataType>
 bool run_grouped_conv_fwd(std::array<ck::index_t, NumDimSpatial + NumNonSpatialDim> in_lengths,
                           std::array<ck::index_t, NumDimSpatial + NumNonSpatialDim> wei_lengths,
                           std::array<ck::index_t, NumDimSpatial + NumNonSpatialDim> out_lengths)

client_example/16_convnd_fwd/conv3d_fwd_fp16_comp_fp8.cpp

@@ -37,7 +37,9 @@ int main()
                                 OutDataType,
                                 InLayout,
                                 WeiLayout,
-                                OutLayout>(
+                                OutLayout,
+                                3,
+                                ck::f8_t>(
                {N, Di, Hi, Wi, G, C}, {G, K, Z, Y, X, C}, {N, Do, Ho, Wo, G, K})
                ? EXIT_SUCCESS
                : EXIT_FAILURE;

library/include/ck/library/tensor_operation_instance/gpu/grouped_convolution_forward.hpp

@@ -378,7 +378,7 @@ template <ck::index_t NumDimSpatial,
           typename InDataType,
           typename WeiDataType,
           typename OutDataType,
-          typename ComputeType = InDataType>
+          typename ComputeType>
 struct DeviceOperationInstanceFactory<ck::tensor_operation::device::DeviceGroupedConvFwdMultipleD<
     NumDimSpatial,
     InLayout,
@@ -521,14 +521,15 @@ struct DeviceOperationInstanceFactory<ck::tensor_operation::device::DeviceGroupe
                 add_device_grouped_conv3d_fwd_xdl_ndhwgc_gkzyxc_ndhwgk_f32_instances(op_ptrs);
             }
             else if constexpr(is_same_v<InDataType, half_t> && is_same_v<WeiDataType, half_t> &&
-                              is_same_v<OutDataType, half_t && is_same_v<ComputeType, half_t>>)
+                              is_same_v<OutDataType, half_t> && is_same_v<ComputeType, half_t>)
             {
                 add_device_grouped_conv3d_fwd_xdl_ndhwgc_gkzyxc_ndhwgk_f16_instances(op_ptrs);
             }
             else if constexpr(is_same_v<InDataType, half_t> && is_same_v<WeiDataType, half_t> &&
-                              is_same_v<OutDataType, half_t && is_same_v<ComputeType, ck::f8_t>>)
+                              is_same_v<OutDataType, half_t> && is_same_v<ComputeType, ck::f8_t>)
             {
-                add_device_grouped_conv3d_fwd_xdl_ndhwgc_gkzyxc_ndhwgk_f16_comp_f8_instances(op_ptrs);
+                add_device_grouped_conv3d_fwd_xdl_ndhwgc_gkzyxc_ndhwgk_f16_comp_f8_instances(
+                    op_ptrs);
             }
             else if constexpr(is_same_v<InDataType, ck::bhalf_t> &&
                               is_same_v<WeiDataType, ck::bhalf_t> &&

library/src/tensor_operation_instance/gpu/grouped_conv3d_fwd/device_grouped_conv3d_fwd_xdl_ndhwgc_gkzyxc_ndhwgk_f16_comp_f8_instance.cpp

@@ -24,27 +24,30 @@ void add_device_grouped_conv3d_fwd_xdl_ndhwgc_gkzyxc_ndhwgk_f16_comp_f8_instance
                                                               PassThrough,
                                                               F8>>>& instances)
 {
-    add_device_operation_instances(instances,
-                                   device_grouped_conv_fwd_xdl_f16_comp_f8_instances<3,
-                                                                             NDHWGC,
-                                                                             GKZYXC,
-                                                                             Empty_Tuple,
-                                                                             NDHWGK,
-                                                                             ConvFwdDefault>{});
-    add_device_operation_instances(instances,
-                                   device_grouped_conv_fwd_xdl_f16_comp_f8_instances<3,
-                                                                             NDHWGC,
-                                                                             GKZYXC,
-                                                                             Empty_Tuple,
-                                                                             NDHWGK,
-                                                                             ConvFwd1x1P0>{});
-    add_device_operation_instances(instances,
-                                   device_grouped_conv_fwd_xdl_f16_comp_f8_instances<3,
-                                                                             NDHWGC,
-                                                                             GKZYXC,
-                                                                             Empty_Tuple,
-                                                                             NDHWGK,
-                                                                             ConvFwd1x1S1P0>{});
+    add_device_operation_instances(
+        instances,
+        device_grouped_conv_fwd_xdl_f16_comp_f8_instances<3,
+                                                          NDHWGC,
+                                                          GKZYXC,
+                                                          Empty_Tuple,
+                                                          NDHWGK,
+                                                          ConvFwdDefault>{});
+    add_device_operation_instances(
+        instances,
+        device_grouped_conv_fwd_xdl_f16_comp_f8_instances<3,
+                                                          NDHWGC,
+                                                          GKZYXC,
+                                                          Empty_Tuple,
+                                                          NDHWGK,
+                                                          ConvFwd1x1P0>{});
+    add_device_operation_instances(
+        instances,
+        device_grouped_conv_fwd_xdl_f16_comp_f8_instances<3,
+                                                          NDHWGC,
+                                                          GKZYXC,
+                                                          Empty_Tuple,
+                                                          NDHWGK,
+                                                          ConvFwd1x1S1P0>{});
 }
 
 } // namespace instance

script/cmake-ck-dev.sh

@@ -6,14 +6,13 @@ rm -rf CMakeFiles
 MY_PROJECT_SOURCE=$1
 
 cmake                                                                                             \
--D INSTANCES_ONLY=ON \
 -D CMAKE_PREFIX_PATH=/opt/rocm                                                                    \
 -D CMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc                                                         \
 -D CMAKE_CXX_FLAGS="-std=c++17 -O3 -ftemplate-backtrace-limit=0  -fPIE  -Wno-gnu-line-marker      \
 -save-temps=$PWD"                                                                                 \
 -D CMAKE_BUILD_TYPE=Release                                                                       \
 -D BUILD_DEV=ON                                                                                   \
--D GPU_TARGETS="gfx908"                                                             \
+-D GPU_TARGETS="gfx908;gfx90a;gfx940"                                                             \
 -D CMAKE_VERBOSE_MAKEFILE:BOOL=ON                                                                 \
 -D USE_BITINT_EXTENSION_INT4=OFF                                                                  \
 ${MY_PROJECT_SOURCE}