Skip to content

Commit

Permalink
Add type related files
Browse files Browse the repository at this point in the history
new file:   include/native_types.hpp
new file:   include/native_types_impl.hpp
new file:   include/type_traits.hpp
new file:   include/types_ext.hpp
new file:   include/xfloat32.hpp
  • Loading branch information
CongMa13 committed Nov 17, 2023
1 parent 082a966 commit b761533
Show file tree
Hide file tree
Showing 23 changed files with 1,734 additions and 91 deletions.
5 changes: 0 additions & 5 deletions library/include/hiptensor/internal/hiptensor_utility.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -131,11 +131,6 @@ namespace std

return os;
}
static ostream& operator<<(ostream& os, const _Float16 value)
{
os << static_cast<float>(value);
return os;
}
}

#endif // HIPTENSOR_UTILITY_INTERNAL_HPP
2 changes: 1 addition & 1 deletion library/src/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -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
)
Expand Down
2 changes: 1 addition & 1 deletion library/src/contraction/contraction_meta_traits.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -34,8 +34,8 @@
#include <element_wise_operation.hpp>

// hiptensor includes
#include "data_types.hpp"
#include "meta_traits.hpp"
#include "types.hpp"

namespace hiptensor
{
Expand Down
2 changes: 1 addition & 1 deletion library/src/contraction/contraction_solution_params.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -30,7 +30,7 @@
#include <hiptensor/hiptensor_types.hpp>

#include "contraction_types.hpp"
#include "types.hpp"
#include "data_types.hpp"

namespace hiptensor
{
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -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
{
Expand Down
2 changes: 1 addition & 1 deletion library/src/contraction/contraction_solution_registry.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -32,8 +32,8 @@
#include <vector>

#include "contraction_types.hpp"
#include "data_types.hpp"
#include "singleton.hpp"
#include "types.hpp"

namespace hiptensor
{
Expand Down
2 changes: 1 addition & 1 deletion library/src/types.cpp → library/src/data_types.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -24,7 +24,7 @@
*
*******************************************************************************/

#include "types.hpp"
#include "data_types.hpp"

namespace hiptensor
{
Expand Down
2 changes: 1 addition & 1 deletion library/src/hiptensor.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
184 changes: 184 additions & 0 deletions library/src/include/config.hpp
Original file line number Diff line number Diff line change
@@ -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
78 changes: 78 additions & 0 deletions library/src/include/data_types.hpp
Original file line number Diff line number Diff line change
@@ -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
Original file line number Diff line number Diff line change
Expand Up @@ -24,10 +24,10 @@
*
*******************************************************************************/

#ifndef HIPTENSOR_LIBRARY_TYPES_IMPL_HPP
#define HIPTENSOR_LIBRARY_TYPES_IMPL_HPP
#ifndef HIPTENSOR_LIBRARY_DATA_TYPES_IMPL_HPP
#define HIPTENSOR_LIBRARY_DATA_TYPES_IMPL_HPP

#include "types.hpp"
#include "data_types.hpp"

namespace hiptensor
{
Expand Down Expand Up @@ -217,4 +217,4 @@ namespace hiptensor

} // namespace hiptensor

#endif // HIPTENSOR_LIBRARY_TYPES_IMPL_HPP
#endif // HIPTENSOR_LIBRARY_DATA_TYPES_IMPL_HPP
Loading

0 comments on commit b761533

Please sign in to comment.