backout some changes

.gitignore

@@ -61,3 +61,7 @@ _doxygen/
 .idea/
 cmake-build*/
 build*/
+
+# Python virtualenv
+.venv/
+

CMakeLists.txt

@@ -9,15 +9,11 @@ get_property(_GENERATOR_IS_MULTI_CONFIG GLOBAL PROPERTY GENERATOR_IS_MULTI_CONFI
 # This has to be initialized before the project() command appears
 # Set the default of CMAKE_BUILD_TYPE to be release, unless user specifies with -D.  MSVC_IDE does not use CMAKE_BUILD_TYPE
 if(_GENERATOR_IS_MULTI_CONFIG)
-    if (NOT CMAKE_CONFIGURATION_TYPES)
-        set(CMAKE_CONFIGURATION_TYPES "Debug;Release;RelWithDebInfo;MinSizeRel" CACHE STRING
-                "Available build types (configurations) on multi-config generators")
-    endif()
+    set(CMAKE_CONFIGURATION_TYPES "Debug;Release;RelWithDebInfo;MinSizeRel" CACHE STRING
+            "Available build types (configurations) on multi-config generators")
 else()
-    if(NOT CMAKE_BUILD_TYPE)
-        set(CMAKE_BUILD_TYPE Release CACHE STRING
-                "Choose the type of build, options are: None Debug Release RelWithDebInfo MinSizeRel.")
-    endif()
+    set(CMAKE_BUILD_TYPE Release CACHE STRING
+            "Choose the type of build, options are: None Debug Release RelWithDebInfo MinSizeRel.")
 endif()
 
 # Default installation path
@@ -171,9 +167,7 @@ if(USE_OPT_NAVI3X)
 endif()
 
 ## Threads
-if(NOT WIN32)
-    set(THREADS_PREFER_PTHREAD_FLAG ON)
-endif()
+set(THREADS_PREFER_PTHREAD_FLAG ON)
 find_package(Threads REQUIRED)
 link_libraries(Threads::Threads)
 

client_example/02_gemm_add_add_fastgelu/CMakeLists.txt

-add_custom_target(client_gemm_fastgelu_examples DEPENDS
-        client_gemm_add_add_fastgelu client_gemm_add_fastgelu client_gemm_fastgelu)
+add_custom_target(client_gemm_fastgelu_examples)
 
 add_executable(client_gemm_add_add_fastgelu gemm_add_add_fastgelu.cpp)
 target_link_libraries(client_gemm_add_add_fastgelu PRIVATE composable_kernel::device_gemm_operations)
@@ -13,8 +12,10 @@ add_executable(client_gemm_fastgelu gemm_fastgelu.cpp)
 target_link_libraries(client_gemm_fastgelu PRIVATE composable_kernel::device_gemm_operations)
 target_compile_features(client_gemm_fastgelu PRIVATE cxx_std_17)
 
-add_custom_target(client_gemm_fastgelu_generic_examples DEPENDS
-        client_gemm_add_add_fastgelu_generic client_gemm_add_fastgelu_generic client_gemm_fastgelu_generic)
+add_dependencies(client_gemm_fastgelu_examples client_gemm_add_add_fastgelu client_gemm_add_fastgelu
+                 client_gemm_fastgelu)
+
+add_custom_target(client_gemm_fastgelu_generic_examples)
 
 add_executable(client_gemm_add_add_fastgelu_generic gemm_add_add_fastgelu_generic.cpp)
 target_link_libraries(client_gemm_add_add_fastgelu_generic composable_kernel::device_gemm_operations)
@@ -27,3 +28,6 @@ target_compile_features(client_gemm_add_fastgelu_generic PRIVATE cxx_std_17)
 add_executable(client_gemm_fastgelu_generic gemm_fastgelu_generic.cpp)
 target_link_libraries(client_gemm_fastgelu_generic PRIVATE composable_kernel::device_gemm_operations)
 target_compile_features(client_gemm_fastgelu_generic PRIVATE cxx_std_17)
+
+add_dependencies(client_gemm_fastgelu_generic_examples client_gemm_add_add_fastgelu_generic 
+                 client_gemm_add_fastgelu_generic client_gemm_fastgelu_generic)
\ No newline at end of file

client_example/05_layernorm/CMakeLists.txt

 add_executable(client_layernorm2d_fwd layernorm2d_fwd.cpp)
 target_link_libraries(client_layernorm2d_fwd PRIVATE composable_kernel::device_other_operations)
-
 target_compile_features(client_layernorm2d_fwd PRIVATE cxx_std_17)
 
 add_executable(client_layernorm4d_fwd layernorm4d_fwd.cpp)
 target_link_libraries(client_layernorm4d_fwd PRIVATE composable_kernel::device_other_operations)
-
 target_compile_features(client_layernorm4d_fwd PRIVATE cxx_std_17)
+

client_example/11_grouped_conv_bwd_weight/CMakeLists.txt

 add_executable(client_grouped_conv1d_bwd_weight_fp16 grouped_conv1d_bwd_weight_fp16.cpp)
-target_link_libraries(client_grouped_conv1d_bwd_weight_fp16 PRIVATE composable_kernel::device_conv_operations)
-target_compile_features(client_grouped_conv1d_bwd_weight_fp16 PRIVATE cxx_std_17)
-
 add_executable(client_grouped_conv2d_bwd_weight_fp16 grouped_conv2d_bwd_weight_fp16.cpp)
-target_link_libraries(client_grouped_conv2d_bwd_weight_fp16 PRIVATE composable_kernel::device_conv_operations)
-target_compile_features(client_grouped_conv2d_bwd_weight_fp16 PRIVATE cxx_std_17)
-
 add_executable(client_grouped_conv3d_bwd_weight_fp16 grouped_conv3d_bwd_weight_fp16.cpp)
-target_link_libraries(client_grouped_conv3d_bwd_weight_fp16 PRIVATE composable_kernel::device_conv_operations)
-target_compile_features(client_grouped_conv3d_bwd_weight_fp16 PRIVATE cxx_std_17)
-
 add_executable(client_grouped_conv3d_bwd_weight_fp32 grouped_conv3d_bwd_weight_fp32.cpp)
-target_link_libraries(client_grouped_conv3d_bwd_weight_fp32 PRIVATE composable_kernel::device_conv_operations)
-target_compile_features(client_grouped_conv3d_bwd_weight_fp32 PRIVATE cxx_std_17)
-
 add_executable(client_grouped_conv3d_bwd_weight_fp16_comp_bf8_fp8 grouped_conv3d_bwd_weight_fp16_comp_bf8_fp8.cpp)
+
+target_link_libraries(client_grouped_conv1d_bwd_weight_fp16 PRIVATE composable_kernel::device_conv_operations)
+target_link_libraries(client_grouped_conv2d_bwd_weight_fp16 PRIVATE composable_kernel::device_conv_operations)
+target_link_libraries(client_grouped_conv3d_bwd_weight_fp16 PRIVATE composable_kernel::device_conv_operations)
+target_link_libraries(client_grouped_conv3d_bwd_weight_fp32 PRIVATE composable_kernel::device_conv_operations)
 target_link_libraries(client_grouped_conv3d_bwd_weight_fp16_comp_bf8_fp8 PRIVATE composable_kernel::device_conv_operations)
+
+target_compile_features(client_grouped_conv1d_bwd_weight_fp16 PRIVATE cxx_std_17)
+target_compile_features(client_grouped_conv2d_bwd_weight_fp16 PRIVATE cxx_std_17)
+target_compile_features(client_grouped_conv3d_bwd_weight_fp16 PRIVATE cxx_std_17)
+target_compile_features(client_grouped_conv3d_bwd_weight_fp32 PRIVATE cxx_std_17)
 target_compile_features(client_grouped_conv3d_bwd_weight_fp16_comp_bf8_fp8 PRIVATE cxx_std_17)

client_example/13_batchnorm/CMakeLists.txt

 add_executable(client_batchnorm_fwd_nhwc batchnorm_fwd_nhwc.cpp)
-target_link_libraries(client_batchnorm_fwd_nhwc PRIVATE composable_kernel::device_other_operations)
-target_compile_features(client_batchnorm_fwd_nhwc PRIVATE cxx_std_17)
-
 add_executable(client_batchnorm_bwd_nhwc batchnorm_bwd_nhwc.cpp)
-target_link_libraries(client_batchnorm_bwd_nhwc PRIVATE composable_kernel::device_other_operations)
-target_compile_features(client_batchnorm_bwd_nhwc PRIVATE cxx_std_17)
-
 add_executable(client_batchnorm_infer_nhwc batchnorm_infer_nhwc.cpp)
+target_link_libraries(client_batchnorm_fwd_nhwc PRIVATE composable_kernel::device_other_operations)
+target_link_libraries(client_batchnorm_bwd_nhwc PRIVATE composable_kernel::device_other_operations)
 target_link_libraries(client_batchnorm_infer_nhwc PRIVATE composable_kernel::device_other_operations)
+target_compile_features(client_batchnorm_fwd_nhwc PRIVATE cxx_std_17)
+target_compile_features(client_batchnorm_bwd_nhwc PRIVATE cxx_std_17)
 target_compile_features(client_batchnorm_infer_nhwc PRIVATE cxx_std_17)

client_example/CMakeLists.txt

@@ -48,31 +48,60 @@ else()
 endif()
 
 find_package(composable_kernel COMPONENTS device_other_operations device_gemm_operations device_conv_operations device_contraction_operations device_reduction_operations)
-find_package(hip REQUIRED PATHS /opt/rocm /opt/rocm/llvm $ENV{HIP_PATH} $ENV{ROCM_PATH})
+find_package(hip REQUIRED PATHS /opt/rocm $ENV{HIP_PATH})
 message(STATUS "Build with HIP ${hip_VERSION}")
 
 add_subdirectory(01_gemm)
-add_subdirectory(02_gemm_add_add_fastgelu)
-add_subdirectory(03_gemm_layernorm)
-add_subdirectory(04_contraction)
-add_subdirectory(05_layernorm)
-add_subdirectory(06_softmax)
-add_subdirectory(07_grouped_convnd_fwd)
-add_subdirectory(08_fused_attention)
-add_subdirectory(09_quantization)
-add_subdirectory(10_grouped_convnd_bwd_data)
-add_subdirectory(11_grouped_conv_bwd_weight)
-add_subdirectory(12_elementwise_normalization)
-add_subdirectory(13_batchnorm)
-add_subdirectory(14_instance_id)
-add_subdirectory(15_convnd_bwd_data)
-add_subdirectory(15_gemm_add_multiply)
-add_subdirectory(15_reduce)
-add_subdirectory(16_convnd_fwd)
-add_subdirectory(17_grouped_gemm_fastgelu)
-add_subdirectory(18_groupnorm)
-add_subdirectory(19_pool)
-add_subdirectory(20_splitk_gemm)
-add_subdirectory(21_grouped_gemm_bias)
-add_subdirectory(22_grouped_gemm)
-add_subdirectory(22_im2col_col2im)
+add_subdirectory(02_gemm_bilinear)
+add_subdirectory(03_gemm_bias_relu)
+add_subdirectory(04_gemm_add_add_fastgelu)
+add_subdirectory(09_convnd_fwd)
+add_subdirectory(10_convnd_fwd_multiple_d_multiple_reduce)
+add_subdirectory(12_reduce)
+add_subdirectory(13_pool2d_fwd)
+add_subdirectory(14_gemm_quantization)
+add_subdirectory(15_grouped_gemm)
+add_subdirectory(16_gemm_multi_d_multi_reduces)
+add_subdirectory(17_convnd_bwd_data)
+add_subdirectory(18_batched_gemm_reduce)
+add_subdirectory(19_binary_elementwise)
+add_subdirectory(20_grouped_conv_bwd_weight)
+add_subdirectory(21_gemm_layernorm)
+add_subdirectory(22_cgemm)
+add_subdirectory(23_softmax)
+add_subdirectory(24_batched_gemm)
+add_subdirectory(25_gemm_bias_e_permute)
+add_subdirectory(26_contraction)
+add_subdirectory(27_layernorm2d_fwd)
+add_subdirectory(28_grouped_gemm_bias_e_permute)
+add_subdirectory(29_batched_gemm_bias_e_permute)
+add_subdirectory(30_grouped_conv_fwd_multiple_d)
+add_subdirectory(31_batched_gemm_gemm)
+add_subdirectory(32_batched_gemm_scale_softmax_gemm)
+add_subdirectory(33_multiple_reduce)
+add_subdirectory(34_batchnorm)
+add_subdirectory(35_splitK_gemm)
+add_subdirectory(36_sparse_embedding)
+add_subdirectory(37_batched_gemm_add_add_relu_gemm_add)
+add_subdirectory(38_grouped_conv_bwd_data_multiple_d)
+add_subdirectory(39_permute)
+add_subdirectory(40_conv2d_fwd_quantization)
+add_subdirectory(41_grouped_conv_conv_fwd)
+add_subdirectory(42_groupnorm_fwd)
+add_subdirectory(43_splitk_gemm_bias_e_permute)
+add_subdirectory(44_elementwise_permute)
+add_subdirectory(45_elementwise_normalization)
+add_subdirectory(46_gemm_add_multiply)
+add_subdirectory(47_gemm_bias_softmax_gemm_permute)
+add_subdirectory(48_pool3d_fwd)
+add_subdirectory(49_maxpool2d_bwd)
+add_subdirectory(50_put_element)
+add_subdirectory(51_avgpool3d_bwd)
+add_subdirectory(52_im2col_col2im)
+add_subdirectory(53_layernorm_bwd)
+add_subdirectory(54_groupnorm_bwd)
+add_subdirectory(60_gemm_multi_ABD)
+add_subdirectory(61_contraction_multi_ABD)
+add_subdirectory(62_conv_fwd_activ)
+add_subdirectory(63_layernorm4d_fwd)
+add_subdirectory(64_tensor_transforms)
\ No newline at end of file

cmake/GTest.cmake

@@ -10,11 +10,13 @@ if(GOOGLETEST_DIR)
 endif()
 
 set(BUILD_GMOCK OFF CACHE INTERNAL "")
+set(INSTALL_GTEST OFF CACHE INTERNAL "")
 
 FetchContent_Declare(
-  googletest
-  GIT_REPOSITORY https://github.com/google/googletest.git
-  GIT_TAG        b85864c64758dec007208e56af933fc3f52044ee
+    googletest
+    GIT_REPOSITORY https://github.com/google/googletest.git
+    GIT_TAG f8d7d77c06936315286eb55f8de22cd23c188571
+    SYSTEM
 )
 
 if(WIN32)
@@ -30,20 +32,6 @@ FetchContent_MakeAvailable(googletest)
 # Restore the old value of BUILD_SHARED_LIBS
 set(BUILD_SHARED_LIBS ${__build_shared_libs} CACHE BOOL "Type of libraries to build" FORCE)
 
-set(GTEST_CMAKE_CXX_FLAGS
-         -Wno-undef
-         -Wno-reserved-identifier
-         -Wno-global-constructors
-         -Wno-missing-noreturn
-         -Wno-disabled-macro-expansion
-         -Wno-used-but-marked-unused
-         -Wno-switch-enum
-         -Wno-zero-as-null-pointer-constant
-         -Wno-unused-member-function
-         -Wno-comma
-         -Wno-old-style-cast
-         -Wno-deprecated
-         -Wno-unsafe-buffer-usage)
 
 if(WIN32)
     list(APPEND GTEST_CMAKE_CXX_FLAGS
@@ -53,12 +41,6 @@ if(WIN32)
             -Wno-language-extension-token)
 endif()
 
-target_compile_options(gtest PRIVATE ${GTEST_CMAKE_CXX_FLAGS})
-target_compile_options(gtest_main PRIVATE ${GTEST_CMAKE_CXX_FLAGS})
-if(WIN32)
-   target_compile_definitions(gtest PUBLIC GTEST_HAS_SEH=0)
-   target_compile_definitions(gtest_main PUBLIC GTEST_HAS_SEH=0)
-endif()
+target_compile_options(gtest PRIVATE -Wno-undef)
+target_compile_options(gtest_main PRIVATE -Wno-undef)
 
-include(GoogleTest)
-unset(GTEST_CMAKE_CXX_FLAGS)

library/src/utility/CMakeLists.txt

@@ -12,7 +12,8 @@ set_target_properties(utility PROPERTIES POSITION_INDEPENDENT_CODE ON)
 target_compile_options(utility PRIVATE ${CMAKE_COMPILER_WARNINGS})
 target_include_directories(utility PUBLIC
     "$<INSTALL_INTERFACE:${CMAKE_INSTALL_INCLUDEDIR}/ck>"
-    "$<INSTALL_INTERFACE:${CMAKE_INSTALL_INCLUDEDIR}/ck/library/utility>")
+    "$<INSTALL_INTERFACE:${CMAKE_INSTALL_INCLUDEDIR}/ck/library/utility>"
+)
 if(WIN32)
     target_compile_definitions(utility PUBLIC NOMINMAX)
 endif()

test/CMakeLists.txt

@@ -6,7 +6,7 @@ include_directories(BEFORE
     ${PROJECT_SOURCE_DIR}/profiler/include
 )
 
-include(GTest)
+include(gtest)
 
 add_custom_target(tests)