Copy wmma type files

.gitignore

@@ -9,7 +9,7 @@ hiptensor_version.hpp
 hiptensor-version.hpp
 
 # Generated source file
-test/01_contraction/configs/*.hpp
+test/*/configs/*.hpp
 
 # Precompiled Headers
 *.gch

CMakeLists.txt

@@ -60,6 +60,7 @@ set(CMAKE_CXX_EXTENSIONS OFF)
 if( CMAKE_PROJECT_NAME STREQUAL "hiptensor" )
   option( HIPTENSOR_BUILD_TESTS "Build hiptensor tests" ON )
   option( HIPTENSOR_BUILD_SAMPLES "Build hiptensor samples" ON )
+  option( HIPTENSOR_DATA_LAYOUT_COL_MAJOR "Set hiptensor data layout to column major" ON )
 endif()
 
 # Setup output paths
@@ -93,6 +94,13 @@ else()
 endif()
 message( VERBOSE "AMDGPU_TARGETS=${AMDGPU_TARGETS}")
 
+if(HIPTENSOR_DATA_LAYOUT_COL_MAJOR)
+  add_compile_definitions(HIPTENSOR_DATA_LAYOUT_COL_MAJOR=1)
+else()
+  add_compile_definitions(HIPTENSOR_DATA_LAYOUT_COL_MAJOR=0)
+endif()
+message("-- HIPTENSOR_DATA_LAYOUT_COL_MAJOR=${HIPTENSOR_DATA_LAYOUT_COL_MAJOR}")
+
 # Setup HIP
 find_package(hip REQUIRED )
 message(STATUS "HIP version: ${hip_VERSION}")

library/src/CMakeLists.txt

@@ -62,7 +62,7 @@ set(HIPTENSOR_CORE_SOURCES
    ${CMAKE_CURRENT_SOURCE_DIR}/hiptensor.cpp
    ${CMAKE_CURRENT_SOURCE_DIR}/logger.cpp
    ${CMAKE_CURRENT_SOURCE_DIR}/performance.cpp
-   ${CMAKE_CURRENT_SOURCE_DIR}/types.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/data_types.cpp
    ${CMAKE_CURRENT_SOURCE_DIR}/hip_device.cpp
    ${CMAKE_CURRENT_SOURCE_DIR}/handle.cpp
 )

library/src/contraction/contraction_meta_traits.hpp

@@ -34,8 +34,8 @@
 #include <element_wise_operation.hpp>
 
 // hiptensor includes
+#include "data_types.hpp"
 #include "meta_traits.hpp"
-#include "types.hpp"
 
 namespace hiptensor
 {

library/src/contraction/contraction_solution_params.hpp

@@ -30,7 +30,7 @@
 #include <hiptensor/hiptensor_types.hpp>
 
 #include "contraction_types.hpp"
-#include "types.hpp"
+#include "data_types.hpp"
 
 namespace hiptensor
 {

library/src/contraction/contraction_solution_params_impl.hpp

@@ -29,8 +29,8 @@
 
 #include "contraction_meta_traits.hpp"
 #include "contraction_solution_params.hpp"
+#include "data_types.hpp"
 #include "hash.hpp"
-#include "types.hpp"
 
 namespace std
 {

library/src/contraction/contraction_solution_registry.hpp

@@ -32,8 +32,8 @@
 #include <vector>
 
 #include "contraction_types.hpp"
+#include "data_types.hpp"
 #include "singleton.hpp"
-#include "types.hpp"
 
 namespace hiptensor
 {

library/src/contraction/hiptensor_contraction.cpp

@@ -87,7 +87,7 @@ hiptensorStatus_t hiptensorInitContractionDescriptor(const hiptensorHandle_t*
     auto& logger = Logger::instance();
 
     // Log API access
-    char msg[1024];
+    char msg[2048];
     snprintf(
         msg,
         sizeof(msg),

library/src/types.cpp → library/src/data_types.cpp

@@ -24,7 +24,7 @@
  *
  *******************************************************************************/
 
-#include "types.hpp"
+#include "data_types.hpp"
 
 namespace hiptensor
 {

library/src/hiptensor.cpp

@@ -27,9 +27,9 @@
 
 #include <hiptensor/hiptensor.hpp>
 
+#include "data_types.hpp"
 #include "handle.hpp"
 #include "logger.hpp"
-#include "types.hpp"
 #include "util.hpp"
 
 hiptensorStatus_t hiptensorCreate(hiptensorHandle_t** handle)
@@ -151,7 +151,8 @@ hiptensorStatus_t hiptensorInitTensorDescriptor(const hiptensorHandle_t*     han
         return HIPTENSOR_STATUS_NOT_INITIALIZED;
     }
 
-    if((lens == nullptr) || ((dataType != HIP_R_32F) && (dataType != HIP_R_64F))
+    if((lens == nullptr)
+       || ((dataType != HIP_R_16F) && (dataType != HIP_R_32F) && (dataType != HIP_R_64F))
        || unaryOp != HIPTENSOR_OP_IDENTITY)
     {
         auto errorCode = HIPTENSOR_STATUS_INVALID_VALUE;

library/src/include/config.hpp

@@ -0,0 +1,184 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2021-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ *
+ *******************************************************************************/
+#ifndef HIPTENSOR_CONFIG_HPP
+#define HIPTENSOR_CONFIG_HPP
+
+namespace hiptensor
+{
+
+///
+/// Architecture support
+/// Guaranteed symbols:
+/// HIPTENSOR_ARCH_GFX908
+/// HIPTENSOR_ARCH_GFX90a
+/// HIPTENSOR_ARCH_GFX940
+/// HIPTENSOR_ARCH_GFX941
+/// HIPTENSOR_ARCH_GFX942
+/// HIPTENSOR_ARCH_GFX1100
+/// HIPTENSOR_ARCH_GFX1101
+/// HIPTENSOR_ARCH_GFX1102
+#if defined(__gfx908__)
+#define HIPTENSOR_ARCH_GFX908 __gfx908__
+#elif defined(__gfx90a__)
+#define HIPTENSOR_ARCH_GFX90A __gfx90a__
+#elif defined(__gfx940__)
+#define HIPTENSOR_ARCH_GFX940 __gfx940__
+#elif defined(__gfx941__)
+#define HIPTENSOR_ARCH_GFX941 __gfx941__
+#elif defined(__gfx942__)
+#define HIPTENSOR_ARCH_GFX942 __gfx942__
+#elif defined(__gfx1100__)
+#define HIPTENSOR_ARCH_GFX1100 __gfx1100__
+#elif defined(__gfx1101__)
+#define HIPTENSOR_ARCH_GFX1101 __gfx1101__
+#elif defined(__gfx1102__)
+#define HIPTENSOR_ARCH_GFX1102 __gfx1102__
+#else
+#define HIPTENSOR_ARCH_HOST 1
+#endif
+
+#if !defined(HIPTENSOR_ARCH_GFX908)
+#define HIPTENSOR_ARCH_GFX908 0
+#endif
+#if !defined(HIPTENSOR_ARCH_GFX90A)
+#define HIPTENSOR_ARCH_GFX90A 0
+#endif
+#if !defined(HIPTENSOR_ARCH_GFX940)
+#define HIPTENSOR_ARCH_GFX940 0
+#endif
+#if !defined(HIPTENSOR_ARCH_GFX941)
+#define HIPTENSOR_ARCH_GFX941 0
+#endif
+#if !defined(HIPTENSOR_ARCH_GFX942)
+#define HIPTENSOR_ARCH_GFX942 0
+#endif
+#if !defined(HIPTENSOR_ARCH_GFX1100)
+#define HIPTENSOR_ARCH_GFX1100 0
+#endif
+#if !defined(HIPTENSOR_ARCH_GFX1101)
+#define HIPTENSOR_ARCH_GFX1101 0
+#endif
+#if !defined(HIPTENSOR_ARCH_GFX1102)
+#define HIPTENSOR_ARCH_GFX1102 0
+#endif
+#if !defined(HIPTENSOR_ARCH_HOST)
+#define HIPTENSOR_ARCH_HOST 0
+#endif
+
+///
+/// Architecture configuration
+/// Guaranteed symbols:
+/// HIPTENSOR_ARCH_GFX9
+/// HIPTENSOR_ARCH_GFX11
+/// HIPTENSOR_WAVE64_MODE
+/// HIPTENSOR_WAVE32_MODE
+/// HIPTENSOR_BLOCK_DIM_16_SUPPORTED
+/// HIPTENSOR_BLOCK_DIM_32_SUPPORTED
+///
+#if HIPTENSOR_ARCH_GFX908 || HIPTENSOR_ARCH_GFX90A || HIPTENSOR_ARCH_GFX940 \
+    || HIPTENSOR_ARCH_GFX941 || HIPTENSOR_ARCH_GFX942
+#define HIPTENSOR_ARCH_GFX9 1
+#define HIPTENSOR_WAVE64_MODE 1
+#define HIPTENSOR_BLOCK_DIM_16_SUPPORTED 1
+#define HIPTENSOR_BLOCK_DIM_32_SUPPORTED 1
+#endif
+
+#if HIPTENSOR_ARCH_GFX1100 || HIPTENSOR_ARCH_GFX1101 || HIPTENSOR_ARCH_GFX1102
+#define HIPTENSOR_ARCH_GFX11 1
+#define HIPTENSOR_WAVE32_MODE 1
+#define HIPTENSOR_BLOCK_DIM_16_SUPPORTED 1
+#endif
+
+#if !defined(HIPTENSOR_ARCH_GFX9)
+#define HIPTENSOR_ARCH_GFX9 0
+#endif
+#if !defined(HIPTENSOR_ARCH_GFX11)
+#define HIPTENSOR_ARCH_GFX11 0
+#endif
+#if !defined(HIPTENSOR_WAVE64_MODE)
+#define HIPTENSOR_WAVE64_MODE 0
+#endif
+#if !defined(HIPTENSOR_WAVE32_MODE)
+#define HIPTENSOR_WAVE32_MODE 0
+#endif
+#if !defined(HIPTENSOR_BLOCK_DIM_16_SUPPORTED)
+#define HIPTENSOR_BLOCK_DIM_16_SUPPORTED 0
+#endif
+#if !defined(HIPTENSOR_BLOCK_DIM_32_SUPPORTED)
+#define HIPTENSOR_BLOCK_DIM_32_SUPPORTED 0
+#endif
+
+#if defined(NDEBUG)
+#define HIPTENSOR_UNSUPPORTED_IMPL(MSG)
+#else
+#define HIPTENSOR_UNSUPPORTED_IMPL(MSG) __attribute__((deprecated(MSG)))
+#endif
+
+#if defined(HIP_NO_HALF)
+#define HIPTENSOR_NO_HALF 1
+#else
+#define HIPTENSOR_NO_HALF 0
+#endif // HIP_NO_HALF
+
+#if HIPTENSOR_NO_HALF || (!HIPTENSOR_NO_HALF && defined(__HIP_NO_HALF_CONVERSIONS__))
+#define HIPTENSOR_TESTS_NO_HALF 1
+#else
+#define HIPTENSOR_TESTS_NO_HALF 0
+#endif // !HIPTENSOR_NO_HALF && defined(__HIP_NO_HALF_CONVERSIONS__)
+
+///
+/// Sanity checks
+///
+#if HIPTENSOR_ARCH_GFX11
+    static_assert((bool)(HIPTENSOR_WAVE32_MODE) && !(bool)(HIPTENSOR_WAVE64_MODE),
+                  "hipTensor supports only wave32 for gfx11 arch");
+    static_assert((bool)(HIPTENSOR_BLOCK_DIM_16_SUPPORTED)
+                      && !(bool)(HIPTENSOR_BLOCK_DIM_32_SUPPORTED),
+                  "hipTensor supports only block size of 16 for gfx11 arch");
+#endif
+
+#if HIPTENSOR_ARCH_GFX9
+    static_assert(!(bool)(HIPTENSOR_WAVE32_MODE) && (bool)(HIPTENSOR_WAVE64_MODE),
+                  "hipTensor supports only wave64 for gfx9 arch");
+    static_assert((bool)(HIPTENSOR_BLOCK_DIM_16_SUPPORTED)
+                      && (bool)(HIPTENSOR_BLOCK_DIM_32_SUPPORTED),
+                  "hipTensor requires block size of 16 and 32 for gfx9 arch");
+#endif
+
+///
+/// Host and Device symbols
+///
+#define HIPTENSOR_DEVICE __device__
+
+#define HIPTENSOR_HOST __host__
+
+#define HIPTENSOR_HOST_DEVICE HIPTENSOR_HOST HIPTENSOR_DEVICE
+
+#define HIPTENSOR_KERNEL __global__
+
+} // namespace hiptensor
+
+#endif // HIPTENSOR_CONFIG_HPP

library/src/include/data_types.hpp

@@ -0,0 +1,78 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ * THE SOFTWARE.
+ *
+ *******************************************************************************/
+
+#ifndef HIPTENSOR_LIBRARY_DATA_TYPES_HPP
+#define HIPTENSOR_LIBRARY_DATA_TYPES_HPP
+
+// clang-format off
+// Include order needs to be preserved
+#include <hip/library_types.h>
+#include <hip/hip_bfloat16.h>
+#include <hip/hip_fp16.h>
+#include <iostream>
+
+#include <hiptensor/hiptensor_types.hpp>
+
+// clang-format on
+
+namespace hiptensor
+{
+    // Used to map to empty tensors
+    struct NoneType;
+
+    static constexpr hipDataType NONE_TYPE = (hipDataType)31;
+
+    // Map type to runtime HipDataType
+    template <typename T>
+    struct HipDataType;
+
+    template <typename T>
+    static constexpr auto HipDataType_v = HipDataType<T>::value;
+
+    // Get data size in bytes from id
+    uint32_t hipDataTypeSize(hipDataType id);
+
+    // Convert hipDataType to hiptensorComputeType_t
+    hiptensorComputeType_t convertToComputeType(hipDataType hipType);
+
+    // Read a single value from void pointer, casted to T
+    template <typename T>
+    T readVal(void const* value, hipDataType id);
+
+    template <typename T>
+    T readVal(void const* value, hiptensorComputeType_t id);
+
+} // namespace hiptensor
+
+bool operator==(hipDataType hipType, hiptensorComputeType_t computeType);
+bool operator==(hiptensorComputeType_t computeType, hipDataType hipType);
+
+bool operator!=(hipDataType hipType, hiptensorComputeType_t computeType);
+bool operator!=(hiptensorComputeType_t computeType, hipDataType hipType);
+
+#include "data_types_impl.hpp"
+
+#endif // HIPTENSOR_LIBRARY_DATA_TYPES_HPP