marsultor/learning-hip
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

should work if we compile it together on the AMD machin in the cloud

Browse Source
This commit is contained in:
Krishna Ayyalasomayajula
2026-01-27 14:28:10 -06:00
commit 536b475d5b
65 changed files with 26551 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

37
.gitignore vendored Normal file
View File

@@ -0,0 +1,37 @@
# OS junk
.DS_Store
Thumbs.db
# Editor settings
.vscode/**
.idea/**
*.swp
*.swo
# Build system metadata (keep actual build dirs tracked)
CMakeFiles/**
CMakeCache.txt
cmake_install.cmake
Makefile
compile_commands.json
# Logs and temp files
*.log
*.tmp
*.bak
*.old
build/**
*.toml
.cache/**
**/*.aux
**/*.fdb_latexmk
**/*.fls
**/*.log
**/*.synctex.gz
**/*.blg
**/*.bcf
**/*.run.xml
**/*.bbl

3
.gitmodules vendored Normal file
View File

@@ -0,0 +1,3 @@
[submodule "external/spdlog"]
path = external/spdlog
url = https://github.com/gabime/spdlog.git

98
CMakeLists.txt Normal file
View File

@@ -0,0 +1,98 @@
cmake_minimum_required(VERSION 3.20)
# -----------------------------
# Force Clang before project()
# -----------------------------
if(NOT DEFINED CMAKE_C_COMPILER)
set(CMAKE_C_COMPILER clang CACHE STRING "" FORCE)
endif()
if(NOT DEFINED CMAKE_CXX_COMPILER)
set(CMAKE_CXX_COMPILER clang++ CACHE STRING "" FORCE)
endif()
# -----------------------------
# Project setup
# -----------------------------
project(learning-hip LANGUAGES CXX)
set(CMAKE_CXX_STANDARD 23)
set(CMAKE_CXX_STANDARD_REQUIRED ON)
set(CMAKE_EXPORT_COMPILE_COMMANDS ON)
set(CMAKE_BUILD_TYPE Release)
# -----------------------------
# System dependency: libsodium
# -----------------------------
find_package(PkgConfig REQUIRED)
# -----------------------------
# Dependencies (vendored)
# -----------------------------
add_subdirectory(external/spdlog)
# -----------------------------
# Executable
# -----------------------------
add_executable(${PROJECT_NAME} src/main.cpp)
# -----------------------------
# Auto-discover include dirs under src/
# -----------------------------
file(GLOB_RECURSE SRC_SUBDIRS LIST_DIRECTORIES true
"${CMAKE_CURRENT_SOURCE_DIR}/src/*"
)
set(SRC_INCLUDES "")
foreach(dir ${SRC_SUBDIRS})
if(IS_DIRECTORY ${dir})
list(APPEND SRC_INCLUDES ${dir})
endif()
endforeach()
# -----------------------------
# Auto-discover all .cpp except main.cpp
# -----------------------------
file(GLOB_RECURSE ALL_CPP
"${CMAKE_CURRENT_SOURCE_DIR}/src/*.cpp"
)
list(REMOVE_ITEM ALL_CPP
"${CMAKE_CURRENT_SOURCE_DIR}/src/main.cpp"
)
# -----------------------------
# Apply sources and includes
# -----------------------------
target_compile_definitions(learning-hip PRIVATE __HIP_PLATFORM_AMD__=1)
target_sources(${PROJECT_NAME} PRIVATE ${ALL_CPP})
target_include_directories(${PROJECT_NAME} PRIVATE
${SRC_INCLUDES}
${CMAKE_CURRENT_SOURCE_DIR}/external/hip # vendored HIP headers
)
# -----------------------------
# Link libraries
# -----------------------------
target_link_libraries(${PROJECT_NAME} PRIVATE
spdlog::spdlog
amdhip64 # HIP runtime only
)
# -----------------------------
# Compiler warnings (Clang / GCC)
# -----------------------------
if(CMAKE_CXX_COMPILER_ID MATCHES "Clang|GNU")
target_compile_options(${PROJECT_NAME} PRIVATE
-Wall
-Wextra
-Wpedantic
-Wconversion
-Wshadow
-Wnull-dereference
-Wdouble-promotion
)
endif()

39
external/hip/channel_descriptor.h vendored Normal file
View File

@@ -0,0 +1,39 @@
/*
Copyright (c) 2015 - 2021 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 HIP_INCLUDE_HIP_CHANNEL_DESCRIPTOR_H
#define HIP_INCLUDE_HIP_CHANNEL_DESCRIPTOR_H
// Some standard header files, these are included by hc.hpp and so want to make them avail on both
// paths to provide a consistent include env and avoid "missing symbol" errors that only appears
// on NVCC path:
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/amd_detail/amd_channel_descriptor.h>
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/nvidia_detail/nvidia_channel_descriptor.h>
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif

38
external/hip/device_functions.h vendored Normal file
View File

@@ -0,0 +1,38 @@
/*
Copyright (c) 2015 - 2023 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 HIP_INCLUDE_HIP_DEVICE_FUNCTIONS_H
#define HIP_INCLUDE_HIP_DEVICE_FUNCTIONS_H
#if !defined(__HIPCC_RTC__)
#include <hip/hip_common.h>
#endif
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/amd_detail/amd_device_functions.h>
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include <device_functions.h>
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif

681
external/hip/driver_types.h vendored Normal file
View File

@@ -0,0 +1,681 @@
/*
Copyright (c) 2015 - 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 HIP_INCLUDE_HIP_DRIVER_TYPES_H
#define HIP_INCLUDE_HIP_DRIVER_TYPES_H
#if !defined(__HIPCC_RTC__)
#include <hip/hip_common.h>
#if __cplusplus
#include <cstdlib>
#else
#include <stdlib.h> // size_t
#endif
#endif
#if !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include "driver_types.h"
#elif defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
/**
* @defgroup DriverTypes Driver Types
* @{
* This section describes the driver data types.
*
*/
typedef void* hipDeviceptr_t;
/**
* HIP channel format kinds
*/
typedef enum hipChannelFormatKind {
hipChannelFormatKindSigned = 0, ///< Signed channel format
hipChannelFormatKindUnsigned = 1, ///< Unsigned channel format
hipChannelFormatKindFloat = 2, ///< Float channel format
hipChannelFormatKindNone = 3 ///< No channel format
} hipChannelFormatKind;
/**
* HIP channel format descriptor
*/
typedef struct hipChannelFormatDesc {
int x;
int y;
int z;
int w;
enum hipChannelFormatKind f; ///< Channel format kind
} hipChannelFormatDesc;
/** @brief The hipTexRefSetArray function flags parameter override format value*/
#define HIP_TRSA_OVERRIDE_FORMAT 0x01
/** @brief The hipTexRefSetFlags function flags parameter read as integer value*/
#define HIP_TRSF_READ_AS_INTEGER 0x01
/** @brief The hipTexRefSetFlags function flags parameter normalized coordinate value*/
#define HIP_TRSF_NORMALIZED_COORDINATES 0x02
/** @brief The hipTexRefSetFlags function flags parameter srgb value*/
#define HIP_TRSF_SRGB 0x10
typedef struct hipArray* hipArray_t;
typedef const struct hipArray* hipArray_const_t;
/**
* HIP array format
*/
typedef enum hipArray_Format {
HIP_AD_FORMAT_UNSIGNED_INT8 = 0x01, ///< Unsigned 8-bit array format
HIP_AD_FORMAT_UNSIGNED_INT16 = 0x02, ///< Unsigned 16-bit array format
HIP_AD_FORMAT_UNSIGNED_INT32 = 0x03, ///< Unsigned 32-bit array format
HIP_AD_FORMAT_SIGNED_INT8 = 0x08, ///< Signed 8-bit array format
HIP_AD_FORMAT_SIGNED_INT16 = 0x09, ///< Signed 16-bit array format
HIP_AD_FORMAT_SIGNED_INT32 = 0x0a, ///< Signed 32-bit array format
HIP_AD_FORMAT_HALF = 0x10, ///< Half array format
HIP_AD_FORMAT_FLOAT = 0x20 ///< Float array format
} hipArray_Format;
/**
* HIP array descriptor
*/
typedef struct HIP_ARRAY_DESCRIPTOR {
size_t Width; ///< Width of the array
size_t Height; ///< Height of the array
enum hipArray_Format Format; ///< Format of the array
unsigned int NumChannels; ///< Number of channels of the array
} HIP_ARRAY_DESCRIPTOR;
/**
* HIP 3D array descriptor
*/
typedef struct HIP_ARRAY3D_DESCRIPTOR {
size_t Width; ///< Width of the array
size_t Height; ///< Height of the array
size_t Depth; ///< Depth of the array
enum hipArray_Format Format; ///< Format of the array
unsigned int NumChannels; ///< Number of channels of the array
unsigned int Flags; ///< Flags of the array
} HIP_ARRAY3D_DESCRIPTOR;
#if !defined(__HIPCC_RTC__)
/**
* HIP 2D memory copy parameters
*/
typedef struct hip_Memcpy2D {
size_t srcXInBytes; ///< Source width in bytes
size_t srcY; ///< Source height
hipMemoryType srcMemoryType; ///< Source memory type
const void* srcHost; ///< Source pointer
hipDeviceptr_t srcDevice; ///< Source device
hipArray_t srcArray; ///< Source array
size_t srcPitch; ///< Source pitch
size_t dstXInBytes; ///< Destination width in bytes
size_t dstY; ///< Destination height
hipMemoryType dstMemoryType; ///< Destination memory type
void* dstHost; ///< Destination pointer
hipDeviceptr_t dstDevice; ///< Destination device
hipArray_t dstArray; ///< Destination array
size_t dstPitch; ///< Destination pitch
size_t WidthInBytes; ///< Width in bytes of the 2D memory copy
size_t Height; ///< Height of the 2D memory copy
} hip_Memcpy2D;
#endif // !defined(__HIPCC_RTC__)
/**
* HIP mipmapped array
*/
typedef struct hipMipmappedArray {
void* data; ///< Data pointer of the mipmapped array
struct hipChannelFormatDesc desc; ///< Description of the mipmapped array
unsigned int type; ///< Type of the mipmapped array
unsigned int width; ///< Width of the mipmapped array
unsigned int height; ///< Height of the mipmapped array
unsigned int depth; ///< Depth of the mipmapped array
unsigned int min_mipmap_level; ///< Minimum level of the mipmapped array
unsigned int max_mipmap_level; ///< Maximum level of the mipmapped array
unsigned int flags; ///< Flags of the mipmapped array
enum hipArray_Format format; ///< Format of the mipmapped array
unsigned int num_channels; ///< Number of channels of the mipmapped array
} hipMipmappedArray;
/**
* HIP mipmapped array pointer
*/
typedef struct hipMipmappedArray* hipMipmappedArray_t;
typedef hipMipmappedArray_t hipmipmappedArray;
typedef const struct hipMipmappedArray* hipMipmappedArray_const_t;
/**
* HIP resource types
*/
typedef enum hipResourceType {
hipResourceTypeArray = 0x00, ///< Array resource
hipResourceTypeMipmappedArray = 0x01, ///< Mipmapped array resource
hipResourceTypeLinear = 0x02, ///< Linear resource
hipResourceTypePitch2D = 0x03 ///< Pitch 2D resource
} hipResourceType;
typedef enum HIPresourcetype_enum {
HIP_RESOURCE_TYPE_ARRAY = 0x00, ///< Array resource
HIP_RESOURCE_TYPE_MIPMAPPED_ARRAY = 0x01, ///< Mipmapped array resource
HIP_RESOURCE_TYPE_LINEAR = 0x02, ///< Linear resource
HIP_RESOURCE_TYPE_PITCH2D = 0x03 ///< Pitch 2D resource
} HIPresourcetype,
hipResourcetype;
/**
* HIP texture address modes
*/
typedef enum HIPaddress_mode_enum {
HIP_TR_ADDRESS_MODE_WRAP = 0, ///< Wrap address mode
HIP_TR_ADDRESS_MODE_CLAMP = 1, ///< Clamp address mode
HIP_TR_ADDRESS_MODE_MIRROR = 2, ///< Mirror address mode
HIP_TR_ADDRESS_MODE_BORDER = 3 ///< Border address mode
} HIPaddress_mode;
/**
* HIP filter modes
*/
typedef enum HIPfilter_mode_enum {
HIP_TR_FILTER_MODE_POINT = 0, ///< Filter mode point
HIP_TR_FILTER_MODE_LINEAR = 1 ///< Filter mode linear
} HIPfilter_mode;
/**
* HIP texture descriptor
*/
typedef struct HIP_TEXTURE_DESC_st {
HIPaddress_mode addressMode[3]; ///< Address modes
HIPfilter_mode filterMode; ///< Filter mode
unsigned int flags; ///< Flags
unsigned int maxAnisotropy; ///< Maximum anisotropy ratio
HIPfilter_mode mipmapFilterMode; ///< Mipmap filter mode
float mipmapLevelBias; ///< Mipmap level bias
float minMipmapLevelClamp; ///< Mipmap minimum level clamp
float maxMipmapLevelClamp; ///< Mipmap maximum level clamp
float borderColor[4]; ///< Border Color
int reserved[12];
} HIP_TEXTURE_DESC;
/**
* HIP texture resource view formats
*/
typedef enum hipResourceViewFormat {
hipResViewFormatNone = 0x00, ///< No resource view format (use underlying resource format)
hipResViewFormatUnsignedChar1 = 0x01, ///< 1 channel, unsigned 8-bit integers
hipResViewFormatUnsignedChar2 = 0x02, ///< 2 channels, unsigned 8-bit integers
hipResViewFormatUnsignedChar4 = 0x03, ///< 4 channels, unsigned 8-bit integers
hipResViewFormatSignedChar1 = 0x04, ///< 1 channel, signed 8-bit integers
hipResViewFormatSignedChar2 = 0x05, ///< 2 channels, signed 8-bit integers
hipResViewFormatSignedChar4 = 0x06, ///< 4 channels, signed 8-bit integers
hipResViewFormatUnsignedShort1 = 0x07, ///< 1 channel, unsigned 16-bit integers
hipResViewFormatUnsignedShort2 = 0x08, ///< 2 channels, unsigned 16-bit integers
hipResViewFormatUnsignedShort4 = 0x09, ///< 4 channels, unsigned 16-bit integers
hipResViewFormatSignedShort1 = 0x0a, ///< 1 channel, signed 16-bit integers
hipResViewFormatSignedShort2 = 0x0b, ///< 2 channels, signed 16-bit integers
hipResViewFormatSignedShort4 = 0x0c, ///< 4 channels, signed 16-bit integers
hipResViewFormatUnsignedInt1 = 0x0d, ///< 1 channel, unsigned 32-bit integers
hipResViewFormatUnsignedInt2 = 0x0e, ///< 2 channels, unsigned 32-bit integers
hipResViewFormatUnsignedInt4 = 0x0f, ///< 4 channels, unsigned 32-bit integers
hipResViewFormatSignedInt1 = 0x10, ///< 1 channel, signed 32-bit integers
hipResViewFormatSignedInt2 = 0x11, ///< 2 channels, signed 32-bit integers
hipResViewFormatSignedInt4 = 0x12, ///< 4 channels, signed 32-bit integers
hipResViewFormatHalf1 = 0x13, ///< 1 channel, 16-bit floating point
hipResViewFormatHalf2 = 0x14, ///< 2 channels, 16-bit floating point
hipResViewFormatHalf4 = 0x15, ///< 4 channels, 16-bit floating point
hipResViewFormatFloat1 = 0x16, ///< 1 channel, 32-bit floating point
hipResViewFormatFloat2 = 0x17, ///< 2 channels, 32-bit floating point
hipResViewFormatFloat4 = 0x18, ///< 4 channels, 32-bit floating point
hipResViewFormatUnsignedBlockCompressed1 = 0x19, ///< Block-compressed 1
hipResViewFormatUnsignedBlockCompressed2 = 0x1a, ///< Block-compressed 2
hipResViewFormatUnsignedBlockCompressed3 = 0x1b, ///< Block-compressed 3
hipResViewFormatUnsignedBlockCompressed4 = 0x1c, ///< Block-compressed 4 unsigned
hipResViewFormatSignedBlockCompressed4 = 0x1d, ///< Block-compressed 4 signed
hipResViewFormatUnsignedBlockCompressed5 = 0x1e, ///< Block-compressed 5 unsigned
hipResViewFormatSignedBlockCompressed5 = 0x1f, ///< Block-compressed 5 signed
hipResViewFormatUnsignedBlockCompressed6H = 0x20, ///< Block-compressed 6 unsigned half-float
hipResViewFormatSignedBlockCompressed6H = 0x21, ///< Block-compressed 6 signed half-float
hipResViewFormatUnsignedBlockCompressed7 = 0x22 ///< Block-compressed 7
} hipResourceViewFormat;
/**
* HIP texture resource view formats
*/
typedef enum HIPresourceViewFormat_enum {
HIP_RES_VIEW_FORMAT_NONE = 0x00, ///< No resource view format (use underlying resource format)
HIP_RES_VIEW_FORMAT_UINT_1X8 = 0x01, ///< 1 channel, unsigned 8-bit integers
HIP_RES_VIEW_FORMAT_UINT_2X8 = 0x02, ///< 2 channels, unsigned 8-bit integers
HIP_RES_VIEW_FORMAT_UINT_4X8 = 0x03, ///< 4 channels, unsigned 8-bit integers
HIP_RES_VIEW_FORMAT_SINT_1X8 = 0x04, ///< 1 channel, signed 8-bit integers
HIP_RES_VIEW_FORMAT_SINT_2X8 = 0x05, ///< 2 channels, signed 8-bit integers
HIP_RES_VIEW_FORMAT_SINT_4X8 = 0x06, ///< 4 channels, signed 8-bit integers
HIP_RES_VIEW_FORMAT_UINT_1X16 = 0x07, ///< 1 channel, unsigned 16-bit integers
HIP_RES_VIEW_FORMAT_UINT_2X16 = 0x08, ///< 2 channels, unsigned 16-bit integers
HIP_RES_VIEW_FORMAT_UINT_4X16 = 0x09, ///< 4 channels, unsigned 16-bit integers
HIP_RES_VIEW_FORMAT_SINT_1X16 = 0x0a, ///< 1 channel, signed 16-bit integers
HIP_RES_VIEW_FORMAT_SINT_2X16 = 0x0b, ///< 2 channels, signed 16-bit integers
HIP_RES_VIEW_FORMAT_SINT_4X16 = 0x0c, ///< 4 channels, signed 16-bit integers
HIP_RES_VIEW_FORMAT_UINT_1X32 = 0x0d, ///< 1 channel, unsigned 32-bit integers
HIP_RES_VIEW_FORMAT_UINT_2X32 = 0x0e, ///< 2 channels, unsigned 32-bit integers
HIP_RES_VIEW_FORMAT_UINT_4X32 = 0x0f, ///< 4 channels, unsigned 32-bit integers
HIP_RES_VIEW_FORMAT_SINT_1X32 = 0x10, ///< 1 channel, signed 32-bit integers
HIP_RES_VIEW_FORMAT_SINT_2X32 = 0x11, ///< 2 channels, signed 32-bit integers
HIP_RES_VIEW_FORMAT_SINT_4X32 = 0x12, ///< 4 channels, signed 32-bit integers
HIP_RES_VIEW_FORMAT_FLOAT_1X16 = 0x13, ///< 1 channel, 16-bit floating point
HIP_RES_VIEW_FORMAT_FLOAT_2X16 = 0x14, ///< 2 channels, 16-bit floating point
HIP_RES_VIEW_FORMAT_FLOAT_4X16 = 0x15, ///< 4 channels, 16-bit floating point
HIP_RES_VIEW_FORMAT_FLOAT_1X32 = 0x16, ///< 1 channel, 32-bit floating point
HIP_RES_VIEW_FORMAT_FLOAT_2X32 = 0x17, ///< 2 channels, 32-bit floating point
HIP_RES_VIEW_FORMAT_FLOAT_4X32 = 0x18, ///< 4 channels, 32-bit floating point
HIP_RES_VIEW_FORMAT_UNSIGNED_BC1 = 0x19, ///< Block-compressed 1
HIP_RES_VIEW_FORMAT_UNSIGNED_BC2 = 0x1a, ///< Block-compressed 2
HIP_RES_VIEW_FORMAT_UNSIGNED_BC3 = 0x1b, ///< Block-compressed 3
HIP_RES_VIEW_FORMAT_UNSIGNED_BC4 = 0x1c, ///< Block-compressed 4 unsigned
HIP_RES_VIEW_FORMAT_SIGNED_BC4 = 0x1d, ///< Block-compressed 4 signed
HIP_RES_VIEW_FORMAT_UNSIGNED_BC5 = 0x1e, ///< Block-compressed 5 unsigned
HIP_RES_VIEW_FORMAT_SIGNED_BC5 = 0x1f, ///< Block-compressed 5 signed
HIP_RES_VIEW_FORMAT_UNSIGNED_BC6H = 0x20, ///< Block-compressed 6 unsigned half-float
HIP_RES_VIEW_FORMAT_SIGNED_BC6H = 0x21, ///< Block-compressed 6 signed half-float
HIP_RES_VIEW_FORMAT_UNSIGNED_BC7 = 0x22 ///< Block-compressed 7
} HIPresourceViewFormat;
/**
* HIP resource descriptor
*/
typedef struct hipResourceDesc {
enum hipResourceType resType; ///< Resource type
union {
struct {
hipArray_t array; ///< HIP array
} array;
struct {
hipMipmappedArray_t mipmap; ///< HIP mipmapped array
} mipmap;
struct {
void* devPtr; ///< Device pointer
struct hipChannelFormatDesc desc; ///< Channel format description
size_t sizeInBytes; ///< Size in bytes
} linear;
struct {
void* devPtr; ///< Device pointer
struct hipChannelFormatDesc desc; ///< Channel format description
size_t width; ///< Width of the array in elements
size_t height; ///< Height of the array in elements
size_t pitchInBytes; ///< Pitch between two rows in bytes
} pitch2D;
} res;
} hipResourceDesc;
/**
* HIP resource view descriptor struct
*/
typedef struct HIP_RESOURCE_DESC_st {
HIPresourcetype resType; ///< Resource type
union {
struct {
hipArray_t hArray; ///< HIP array
} array;
struct {
hipMipmappedArray_t hMipmappedArray; ///< HIP mipmapped array
} mipmap;
struct {
hipDeviceptr_t devPtr; ///< Device pointer
hipArray_Format format; ///< Array format
unsigned int numChannels; ///< Channels per array element
size_t sizeInBytes; ///< Size in bytes
} linear;
struct {
hipDeviceptr_t devPtr; ///< Device pointer
hipArray_Format format; ///< Array format
unsigned int numChannels; ///< Channels per array element
size_t width; ///< Width of the array in elements
size_t height; ///< Height of the array in elements
size_t pitchInBytes; ///< Pitch between two rows in bytes
} pitch2D;
struct {
int reserved[32];
} reserved;
} res;
unsigned int flags; ///< Flags (must be zero)
} HIP_RESOURCE_DESC;
/**
* HIP resource view descriptor
*/
struct hipResourceViewDesc {
enum hipResourceViewFormat format; ///< Resource view format
size_t width; ///< Width of the resource view
size_t height; ///< Height of the resource view
size_t depth; ///< Depth of the resource view
unsigned int firstMipmapLevel; ///< First defined mipmap level
unsigned int lastMipmapLevel; ///< Last defined mipmap level
unsigned int firstLayer; ///< First layer index
unsigned int lastLayer; ///< Last layer index
};
/**
* Resource view descriptor
*/
typedef struct HIP_RESOURCE_VIEW_DESC_st {
HIPresourceViewFormat format; ///< Resource view format
size_t width; ///< Width of the resource view
size_t height; ///< Height of the resource view
size_t depth; ///< Depth of the resource view
unsigned int firstMipmapLevel; ///< First defined mipmap level
unsigned int lastMipmapLevel; ///< Last defined mipmap level
unsigned int firstLayer; ///< First layer index
unsigned int lastLayer; ///< Last layer index
unsigned int reserved[16];
} HIP_RESOURCE_VIEW_DESC;
/**
* Memory copy types
*/
#if !defined(__HIPCC_RTC__)
typedef enum hipMemcpyKind {
hipMemcpyHostToHost = 0, ///< Host-to-Host Copy
hipMemcpyHostToDevice = 1, ///< Host-to-Device Copy
hipMemcpyDeviceToHost = 2, ///< Device-to-Host Copy
hipMemcpyDeviceToDevice = 3, ///< Device-to-Device Copy
hipMemcpyDefault = 4, ///< Runtime will automatically determine
///< copy-kind based on virtual addresses.
hipMemcpyDeviceToDeviceNoCU = 1024 ///< Device-to-Device Copy without using compute units
} hipMemcpyKind;
/**
* HIP pithed pointer
*/
typedef struct hipPitchedPtr {
void* ptr; ///< Pointer to the allocated memory
size_t pitch; ///< Pitch in bytes
size_t xsize; ///< Logical size of the first dimension of allocation in elements
size_t ysize; ///< Logical size of the second dimension of allocation in elements
} hipPitchedPtr;
/**
* HIP extent
*/
typedef struct hipExtent {
size_t width; // Width in elements when referring to array memory, in bytes when referring to
// linear memory
size_t height;
size_t depth;
} hipExtent;
/**
* HIP position
*/
typedef struct hipPos {
size_t x; ///< X coordinate
size_t y; ///< Y coordinate
size_t z; ///< Z coordinate
} hipPos;
/**
* HIP 3D memory copy parameters
*/
typedef struct hipMemcpy3DParms {
hipArray_t srcArray; ///< Source array
struct hipPos srcPos; ///< Source position
struct hipPitchedPtr srcPtr; ///< Source pointer
hipArray_t dstArray; ///< Destination array
struct hipPos dstPos; ///< Destination position
struct hipPitchedPtr dstPtr; ///< Destination pointer
struct hipExtent extent; ///< Extent of 3D memory copy
enum hipMemcpyKind kind; ///< Kind of 3D memory copy
} hipMemcpy3DParms;
/**
* HIP 3D memory copy
*/
typedef struct HIP_MEMCPY3D {
size_t srcXInBytes; ///< Source X in bytes
size_t srcY; ///< Source Y
size_t srcZ; ///< Source Z
size_t srcLOD; ///< Source LOD
hipMemoryType srcMemoryType; ///< Source memory type
const void* srcHost; ///< Source host pointer
hipDeviceptr_t srcDevice; ///< Source device
hipArray_t srcArray; ///< Source array
size_t srcPitch; ///< Source pitch
size_t srcHeight; ///< Source height
size_t dstXInBytes; ///< Destination X in bytes
size_t dstY; ///< Destination Y
size_t dstZ; ///< Destination Z
size_t dstLOD; ///< Destination LOD
hipMemoryType dstMemoryType; ///< Destination memory type
void* dstHost; ///< Destination host pointer
hipDeviceptr_t dstDevice; ///< Destination device
hipArray_t dstArray; ///< Destination array
size_t dstPitch; ///< Destination pitch
size_t dstHeight; ///< Destination height
size_t WidthInBytes; ///< Width in bytes of 3D memory copy
size_t Height; ///< Height in bytes of 3D memory copy
size_t Depth; ///< Depth in bytes of 3D memory copy
} HIP_MEMCPY3D;
/**
* Specifies the type of location
*/
typedef enum hipMemLocationType {
hipMemLocationTypeInvalid = 0,
hipMemLocationTypeNone = 0,
hipMemLocationTypeDevice = 1, ///< Device location, thus it's HIP device ID
hipMemLocationTypeHost = 2, ///< Host location, id is ignored
hipMemLocationTypeHostNuma = 3, ///< Host NUMA node location, id is host NUMA node id
hipMemLocationTypeHostNumaCurrent =
4 ///< Host NUMA node closest to current thread’s CPU, id is ignored
} hipMemLocationType;
/**
* Specifies a memory location.
*
* To specify a gpu, set type = @p hipMemLocationTypeDevice and set id = the gpu's device ID
*/
typedef struct hipMemLocation {
hipMemLocationType type; ///< Specifies the location type, which describes the meaning of id
int id; ///< Identifier for the provided location type @p hipMemLocationType
} hipMemLocation;
/**
* Flags to specify for copies within a batch. Used with hipMemcpyBatchAsync
*/
typedef enum hipMemcpyFlags {
hipMemcpyFlagDefault = 0x0, ///< Default flag
hipMemcpyFlagPreferOverlapWithCompute = 0x1 ///< Tries to overlap copy with compute work.
} hipMemcpyFlags;
/**
* Flags to specify order in which source pointer is accessed by Batch memcpy
*/
typedef enum hipMemcpySrcAccessOrder {
hipMemcpySrcAccessOrderInvalid = 0x0, ///< Default Invalid.
hipMemcpySrcAccessOrderStream = 0x1, ///< Access to source pointer must be in stream order.
hipMemcpySrcAccessOrderDuringApiCall =
0x2, ///< Access to source pointer can be out of stream order and all accesses must be
///< complete before API call returns.
hipMemcpySrcAccessOrderAny =
0x3, ///< Access to the source pointer can be out of stream order and the accesses can happen
///< even after the API call return.
hipMemcpySrcAccessOrderMax = 0x7FFFFFFF
} hipMemcpySrcAccessOrder;
/**
* Attributes for copies within a batch.
*/
typedef struct hipMemcpyAttributes {
hipMemcpySrcAccessOrder
srcAccessOrder; ///< Source access ordering to be observed for copies with this attribute.
hipMemLocation srcLocHint; ///< Location hint for src operand.
hipMemLocation dstLocHint; ///< Location hint for destination operand.
unsigned int flags; ///< Additional Flags for copies. See hipMemcpyFlags.
} hipMemcpyAttributes;
/**
* Operand types for individual copies within a batch
*/
typedef enum hipMemcpy3DOperandType {
hipMemcpyOperandTypePointer = 0x1, ///< Mempcy operand is a valid pointer.
hipMemcpyOperandTypeArray = 0x2, ///< Memcpy operand is a valid hipArray.
hipMemcpyOperandTypeMax = 0x7FFFFFFF
} hipMemcpy3DOperandType;
/**
* Struct representing offset into a hipArray_t in elements.
*/
typedef struct hipOffset3D {
size_t x;
size_t y;
size_t z;
} hipOffset3D;
/**
* Struct representing an operand for copy with hipMemcpy3DBatchAsync.
*/
typedef struct hipMemcpy3DOperand {
hipMemcpy3DOperandType type;
union {
struct {
void* ptr;
size_t rowLength; ///< Length of each row in elements.
size_t layerHeight; ///< Height of each layer in elements.
hipMemLocation locHint; ///< Location Hint for the operand.
} ptr;
struct {
hipArray_t array; ///< Array struct for hipMemcpyOperandTypeArray.
hipOffset3D offset; ///< Offset into array in elements.
} array;
} op;
} hipMemcpy3DOperand;
/**
* HIP 3D Batch Op
*/
typedef struct hipMemcpy3DBatchOp {
hipMemcpy3DOperand src;
hipMemcpy3DOperand dst;
hipExtent extent;
hipMemcpySrcAccessOrder srcAccessOrder;
unsigned int flags;
} hipMemcpy3DBatchOp;
typedef struct hipMemcpy3DPeerParms {
hipArray_t srcArray; ///< Source memory address
hipPos srcPos; ///< Source position offset
hipPitchedPtr srcPtr; ///< Pitched source memory address
int srcDevice; ///< Source device
hipArray_t dstArray; ///< Destination memory address
hipPos dstPos; ///< Destination position offset
hipPitchedPtr dstPtr; ///< Pitched destination memory address
int dstDevice; ///< Destination device
hipExtent extent; ///< Requested memory copy size
} hipMemcpy3DPeerParms;
/**
* @brief Make hipPitchedPtr
*
* @param [in] d Pointer to the allocated memory
* @param [in] p Pitch in bytes
* @param [in] xsz Logical size of the first dimension of allocation in elements
* @param [in] ysz Logical size of the second dimension of allocation in elements
*
* @returns The created hipPitchedPtr
*/
static inline struct hipPitchedPtr make_hipPitchedPtr(void* d, size_t p, size_t xsz, size_t ysz) {
struct hipPitchedPtr s;
s.ptr = d;
s.pitch = p;
s.xsize = xsz;
s.ysize = ysz;
return s;
}
/**
* @brief Make hipPos struct
*
* @param [in] x X coordinate of the new hipPos
* @param [in] y Y coordinate of the new hipPos
* @param [in] z Z coordinate of the new hipPos
*
* @returns The created hipPos struct
*/
static inline struct hipPos make_hipPos(size_t x, size_t y, size_t z) {
struct hipPos p;
p.x = x;
p.y = y;
p.z = z;
return p;
}
/**
* @brief Make hipExtent struct
*
* @param [in] w Width of the new hipExtent
* @param [in] h Height of the new hipExtent
* @param [in] d Depth of the new hipExtent
*
* @returns The created hipExtent struct
*/
static inline struct hipExtent make_hipExtent(size_t w, size_t h, size_t d) {
struct hipExtent e;
e.width = w;
e.height = h;
e.depth = d;
return e;
}
typedef enum hipFunction_attribute {
HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, ///< The maximum number of threads per block. Depends
///< on function and device.
HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES, ///< The statically allocated shared memory size in bytes
///< per block required by the function.
HIP_FUNC_ATTRIBUTE_CONST_SIZE_BYTES, ///< The user-allocated constant memory by the function in
///< bytes.
HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES, ///< The local memory usage of each thread by this function
///< in bytes.
HIP_FUNC_ATTRIBUTE_NUM_REGS, ///< The number of registers used by each thread of this function.
HIP_FUNC_ATTRIBUTE_PTX_VERSION, ///< PTX version
HIP_FUNC_ATTRIBUTE_BINARY_VERSION, ///< Binary version
HIP_FUNC_ATTRIBUTE_CACHE_MODE_CA, ///< Cache mode
HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, ///< The maximum dynamic shared memory per
///< block for this function in bytes.
HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT, ///< The shared memory carveout preference
///< in percent of the maximum shared
///< memory.
HIP_FUNC_ATTRIBUTE_MAX
} hipFunction_attribute;
typedef enum hipPointer_attribute {
HIP_POINTER_ATTRIBUTE_CONTEXT = 1, ///< The context on which a pointer was allocated
///< @warning This attribute is not supported in HIP
HIP_POINTER_ATTRIBUTE_MEMORY_TYPE, ///< memory type describing the location of a pointer
HIP_POINTER_ATTRIBUTE_DEVICE_POINTER, ///< address at which the pointer is allocated on the
///< device
HIP_POINTER_ATTRIBUTE_HOST_POINTER, ///< address at which the pointer is allocated on the host
HIP_POINTER_ATTRIBUTE_P2P_TOKENS, ///< A pair of tokens for use with Linux kernel interface
///< @warning This attribute is not supported in HIP
HIP_POINTER_ATTRIBUTE_SYNC_MEMOPS, ///< Synchronize every synchronous memory operation
///< initiated on this region
HIP_POINTER_ATTRIBUTE_BUFFER_ID, ///< Unique ID for an allocated memory region
HIP_POINTER_ATTRIBUTE_IS_MANAGED, ///< Indicates if the pointer points to managed memory
HIP_POINTER_ATTRIBUTE_DEVICE_ORDINAL, ///< device ordinal of a device on which a pointer
///< was allocated or registered
HIP_POINTER_ATTRIBUTE_IS_LEGACY_HIP_IPC_CAPABLE, ///< if this pointer maps to an allocation
///< that is suitable for hipIpcGetMemHandle
///< @warning This attribute is not supported in
///< HIP
HIP_POINTER_ATTRIBUTE_RANGE_START_ADDR, ///< Starting address for this requested pointer
HIP_POINTER_ATTRIBUTE_RANGE_SIZE, ///< Size of the address range for this requested pointer
HIP_POINTER_ATTRIBUTE_MAPPED, ///< tells if this pointer is in a valid address range
///< that is mapped to a backing allocation
HIP_POINTER_ATTRIBUTE_ALLOWED_HANDLE_TYPES, ///< Bitmask of allowed hipmemAllocationHandleType
///< for this allocation @warning This attribute is
///< not supported in HIP
HIP_POINTER_ATTRIBUTE_IS_GPU_DIRECT_RDMA_CAPABLE, ///< returns if the memory referenced by
///< this pointer can be used with the
///< GPUDirect RDMA API
///< @warning This attribute is not supported
///< in HIP
HIP_POINTER_ATTRIBUTE_ACCESS_FLAGS, ///< Returns the access flags the device associated with
///< for the corresponding memory referenced by the ptr
HIP_POINTER_ATTRIBUTE_MEMPOOL_HANDLE ///< Returns the mempool handle for the allocation if
///< it was allocated from a mempool
///< @warning This attribute is not supported in HIP
} hipPointer_attribute;
// doxygen end DriverTypes
/**
* @}
*/
#endif // !defined(__HIPCC_RTC__)
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif

39
external/hip/hip/channel_descriptor.h vendored Normal file
View File

@@ -0,0 +1,39 @@
/*
Copyright (c) 2015 - 2021 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 HIP_INCLUDE_HIP_CHANNEL_DESCRIPTOR_H
#define HIP_INCLUDE_HIP_CHANNEL_DESCRIPTOR_H
// Some standard header files, these are included by hc.hpp and so want to make them avail on both
// paths to provide a consistent include env and avoid "missing symbol" errors that only appears
// on NVCC path:
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/amd_detail/amd_channel_descriptor.h>
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/nvidia_detail/nvidia_channel_descriptor.h>
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif

38
external/hip/hip/device_functions.h vendored Normal file
View File

@@ -0,0 +1,38 @@
/*
Copyright (c) 2015 - 2023 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 HIP_INCLUDE_HIP_DEVICE_FUNCTIONS_H
#define HIP_INCLUDE_HIP_DEVICE_FUNCTIONS_H
#if !defined(__HIPCC_RTC__)
#include <hip/hip_common.h>
#endif
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/amd_detail/amd_device_functions.h>
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include <device_functions.h>
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif

681
external/hip/hip/driver_types.h vendored Normal file
View File

@@ -0,0 +1,681 @@
/*
Copyright (c) 2015 - 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 HIP_INCLUDE_HIP_DRIVER_TYPES_H
#define HIP_INCLUDE_HIP_DRIVER_TYPES_H
#if !defined(__HIPCC_RTC__)
#include <hip/hip_common.h>
#if __cplusplus
#include <cstdlib>
#else
#include <stdlib.h> // size_t
#endif
#endif
#if !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include "driver_types.h"
#elif defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
/**
* @defgroup DriverTypes Driver Types
* @{
* This section describes the driver data types.
*
*/
typedef void* hipDeviceptr_t;
/**
* HIP channel format kinds
*/
typedef enum hipChannelFormatKind {
hipChannelFormatKindSigned = 0, ///< Signed channel format
hipChannelFormatKindUnsigned = 1, ///< Unsigned channel format
hipChannelFormatKindFloat = 2, ///< Float channel format
hipChannelFormatKindNone = 3 ///< No channel format
} hipChannelFormatKind;
/**
* HIP channel format descriptor
*/
typedef struct hipChannelFormatDesc {
int x;
int y;
int z;
int w;
enum hipChannelFormatKind f; ///< Channel format kind
} hipChannelFormatDesc;
/** @brief The hipTexRefSetArray function flags parameter override format value*/
#define HIP_TRSA_OVERRIDE_FORMAT 0x01
/** @brief The hipTexRefSetFlags function flags parameter read as integer value*/
#define HIP_TRSF_READ_AS_INTEGER 0x01
/** @brief The hipTexRefSetFlags function flags parameter normalized coordinate value*/
#define HIP_TRSF_NORMALIZED_COORDINATES 0x02
/** @brief The hipTexRefSetFlags function flags parameter srgb value*/
#define HIP_TRSF_SRGB 0x10
typedef struct hipArray* hipArray_t;
typedef const struct hipArray* hipArray_const_t;
/**
* HIP array format
*/
typedef enum hipArray_Format {
HIP_AD_FORMAT_UNSIGNED_INT8 = 0x01, ///< Unsigned 8-bit array format
HIP_AD_FORMAT_UNSIGNED_INT16 = 0x02, ///< Unsigned 16-bit array format
HIP_AD_FORMAT_UNSIGNED_INT32 = 0x03, ///< Unsigned 32-bit array format
HIP_AD_FORMAT_SIGNED_INT8 = 0x08, ///< Signed 8-bit array format
HIP_AD_FORMAT_SIGNED_INT16 = 0x09, ///< Signed 16-bit array format
HIP_AD_FORMAT_SIGNED_INT32 = 0x0a, ///< Signed 32-bit array format
HIP_AD_FORMAT_HALF = 0x10, ///< Half array format
HIP_AD_FORMAT_FLOAT = 0x20 ///< Float array format
} hipArray_Format;
/**
* HIP array descriptor
*/
typedef struct HIP_ARRAY_DESCRIPTOR {
size_t Width; ///< Width of the array
size_t Height; ///< Height of the array
enum hipArray_Format Format; ///< Format of the array
unsigned int NumChannels; ///< Number of channels of the array
} HIP_ARRAY_DESCRIPTOR;
/**
* HIP 3D array descriptor
*/
typedef struct HIP_ARRAY3D_DESCRIPTOR {
size_t Width; ///< Width of the array
size_t Height; ///< Height of the array
size_t Depth; ///< Depth of the array
enum hipArray_Format Format; ///< Format of the array
unsigned int NumChannels; ///< Number of channels of the array
unsigned int Flags; ///< Flags of the array
} HIP_ARRAY3D_DESCRIPTOR;
#if !defined(__HIPCC_RTC__)
/**
* HIP 2D memory copy parameters
*/
typedef struct hip_Memcpy2D {
size_t srcXInBytes; ///< Source width in bytes
size_t srcY; ///< Source height
hipMemoryType srcMemoryType; ///< Source memory type
const void* srcHost; ///< Source pointer
hipDeviceptr_t srcDevice; ///< Source device
hipArray_t srcArray; ///< Source array
size_t srcPitch; ///< Source pitch
size_t dstXInBytes; ///< Destination width in bytes
size_t dstY; ///< Destination height
hipMemoryType dstMemoryType; ///< Destination memory type
void* dstHost; ///< Destination pointer
hipDeviceptr_t dstDevice; ///< Destination device
hipArray_t dstArray; ///< Destination array
size_t dstPitch; ///< Destination pitch
size_t WidthInBytes; ///< Width in bytes of the 2D memory copy
size_t Height; ///< Height of the 2D memory copy
} hip_Memcpy2D;
#endif // !defined(__HIPCC_RTC__)
/**
* HIP mipmapped array
*/
typedef struct hipMipmappedArray {
void* data; ///< Data pointer of the mipmapped array
struct hipChannelFormatDesc desc; ///< Description of the mipmapped array
unsigned int type; ///< Type of the mipmapped array
unsigned int width; ///< Width of the mipmapped array
unsigned int height; ///< Height of the mipmapped array
unsigned int depth; ///< Depth of the mipmapped array
unsigned int min_mipmap_level; ///< Minimum level of the mipmapped array
unsigned int max_mipmap_level; ///< Maximum level of the mipmapped array
unsigned int flags; ///< Flags of the mipmapped array
enum hipArray_Format format; ///< Format of the mipmapped array
unsigned int num_channels; ///< Number of channels of the mipmapped array
} hipMipmappedArray;
/**
* HIP mipmapped array pointer
*/
typedef struct hipMipmappedArray* hipMipmappedArray_t;
typedef hipMipmappedArray_t hipmipmappedArray;
typedef const struct hipMipmappedArray* hipMipmappedArray_const_t;
/**
* HIP resource types
*/
typedef enum hipResourceType {
hipResourceTypeArray = 0x00, ///< Array resource
hipResourceTypeMipmappedArray = 0x01, ///< Mipmapped array resource
hipResourceTypeLinear = 0x02, ///< Linear resource
hipResourceTypePitch2D = 0x03 ///< Pitch 2D resource
} hipResourceType;
typedef enum HIPresourcetype_enum {
HIP_RESOURCE_TYPE_ARRAY = 0x00, ///< Array resource
HIP_RESOURCE_TYPE_MIPMAPPED_ARRAY = 0x01, ///< Mipmapped array resource
HIP_RESOURCE_TYPE_LINEAR = 0x02, ///< Linear resource
HIP_RESOURCE_TYPE_PITCH2D = 0x03 ///< Pitch 2D resource
} HIPresourcetype,
hipResourcetype;
/**
* HIP texture address modes
*/
typedef enum HIPaddress_mode_enum {
HIP_TR_ADDRESS_MODE_WRAP = 0, ///< Wrap address mode
HIP_TR_ADDRESS_MODE_CLAMP = 1, ///< Clamp address mode
HIP_TR_ADDRESS_MODE_MIRROR = 2, ///< Mirror address mode
HIP_TR_ADDRESS_MODE_BORDER = 3 ///< Border address mode
} HIPaddress_mode;
/**
* HIP filter modes
*/
typedef enum HIPfilter_mode_enum {
HIP_TR_FILTER_MODE_POINT = 0, ///< Filter mode point
HIP_TR_FILTER_MODE_LINEAR = 1 ///< Filter mode linear
} HIPfilter_mode;
/**
* HIP texture descriptor
*/
typedef struct HIP_TEXTURE_DESC_st {
HIPaddress_mode addressMode[3]; ///< Address modes
HIPfilter_mode filterMode; ///< Filter mode
unsigned int flags; ///< Flags
unsigned int maxAnisotropy; ///< Maximum anisotropy ratio
HIPfilter_mode mipmapFilterMode; ///< Mipmap filter mode
float mipmapLevelBias; ///< Mipmap level bias
float minMipmapLevelClamp; ///< Mipmap minimum level clamp
float maxMipmapLevelClamp; ///< Mipmap maximum level clamp
float borderColor[4]; ///< Border Color
int reserved[12];
} HIP_TEXTURE_DESC;
/**
* HIP texture resource view formats
*/
typedef enum hipResourceViewFormat {
hipResViewFormatNone = 0x00, ///< No resource view format (use underlying resource format)
hipResViewFormatUnsignedChar1 = 0x01, ///< 1 channel, unsigned 8-bit integers
hipResViewFormatUnsignedChar2 = 0x02, ///< 2 channels, unsigned 8-bit integers
hipResViewFormatUnsignedChar4 = 0x03, ///< 4 channels, unsigned 8-bit integers
hipResViewFormatSignedChar1 = 0x04, ///< 1 channel, signed 8-bit integers
hipResViewFormatSignedChar2 = 0x05, ///< 2 channels, signed 8-bit integers
hipResViewFormatSignedChar4 = 0x06, ///< 4 channels, signed 8-bit integers
hipResViewFormatUnsignedShort1 = 0x07, ///< 1 channel, unsigned 16-bit integers
hipResViewFormatUnsignedShort2 = 0x08, ///< 2 channels, unsigned 16-bit integers
hipResViewFormatUnsignedShort4 = 0x09, ///< 4 channels, unsigned 16-bit integers
hipResViewFormatSignedShort1 = 0x0a, ///< 1 channel, signed 16-bit integers
hipResViewFormatSignedShort2 = 0x0b, ///< 2 channels, signed 16-bit integers
hipResViewFormatSignedShort4 = 0x0c, ///< 4 channels, signed 16-bit integers
hipResViewFormatUnsignedInt1 = 0x0d, ///< 1 channel, unsigned 32-bit integers
hipResViewFormatUnsignedInt2 = 0x0e, ///< 2 channels, unsigned 32-bit integers
hipResViewFormatUnsignedInt4 = 0x0f, ///< 4 channels, unsigned 32-bit integers
hipResViewFormatSignedInt1 = 0x10, ///< 1 channel, signed 32-bit integers
hipResViewFormatSignedInt2 = 0x11, ///< 2 channels, signed 32-bit integers
hipResViewFormatSignedInt4 = 0x12, ///< 4 channels, signed 32-bit integers
hipResViewFormatHalf1 = 0x13, ///< 1 channel, 16-bit floating point
hipResViewFormatHalf2 = 0x14, ///< 2 channels, 16-bit floating point
hipResViewFormatHalf4 = 0x15, ///< 4 channels, 16-bit floating point
hipResViewFormatFloat1 = 0x16, ///< 1 channel, 32-bit floating point
hipResViewFormatFloat2 = 0x17, ///< 2 channels, 32-bit floating point
hipResViewFormatFloat4 = 0x18, ///< 4 channels, 32-bit floating point
hipResViewFormatUnsignedBlockCompressed1 = 0x19, ///< Block-compressed 1
hipResViewFormatUnsignedBlockCompressed2 = 0x1a, ///< Block-compressed 2
hipResViewFormatUnsignedBlockCompressed3 = 0x1b, ///< Block-compressed 3
hipResViewFormatUnsignedBlockCompressed4 = 0x1c, ///< Block-compressed 4 unsigned
hipResViewFormatSignedBlockCompressed4 = 0x1d, ///< Block-compressed 4 signed
hipResViewFormatUnsignedBlockCompressed5 = 0x1e, ///< Block-compressed 5 unsigned
hipResViewFormatSignedBlockCompressed5 = 0x1f, ///< Block-compressed 5 signed
hipResViewFormatUnsignedBlockCompressed6H = 0x20, ///< Block-compressed 6 unsigned half-float
hipResViewFormatSignedBlockCompressed6H = 0x21, ///< Block-compressed 6 signed half-float
hipResViewFormatUnsignedBlockCompressed7 = 0x22 ///< Block-compressed 7
} hipResourceViewFormat;
/**
* HIP texture resource view formats
*/
typedef enum HIPresourceViewFormat_enum {
HIP_RES_VIEW_FORMAT_NONE = 0x00, ///< No resource view format (use underlying resource format)
HIP_RES_VIEW_FORMAT_UINT_1X8 = 0x01, ///< 1 channel, unsigned 8-bit integers
HIP_RES_VIEW_FORMAT_UINT_2X8 = 0x02, ///< 2 channels, unsigned 8-bit integers
HIP_RES_VIEW_FORMAT_UINT_4X8 = 0x03, ///< 4 channels, unsigned 8-bit integers
HIP_RES_VIEW_FORMAT_SINT_1X8 = 0x04, ///< 1 channel, signed 8-bit integers
HIP_RES_VIEW_FORMAT_SINT_2X8 = 0x05, ///< 2 channels, signed 8-bit integers
HIP_RES_VIEW_FORMAT_SINT_4X8 = 0x06, ///< 4 channels, signed 8-bit integers
HIP_RES_VIEW_FORMAT_UINT_1X16 = 0x07, ///< 1 channel, unsigned 16-bit integers
HIP_RES_VIEW_FORMAT_UINT_2X16 = 0x08, ///< 2 channels, unsigned 16-bit integers
HIP_RES_VIEW_FORMAT_UINT_4X16 = 0x09, ///< 4 channels, unsigned 16-bit integers
HIP_RES_VIEW_FORMAT_SINT_1X16 = 0x0a, ///< 1 channel, signed 16-bit integers
HIP_RES_VIEW_FORMAT_SINT_2X16 = 0x0b, ///< 2 channels, signed 16-bit integers
HIP_RES_VIEW_FORMAT_SINT_4X16 = 0x0c, ///< 4 channels, signed 16-bit integers
HIP_RES_VIEW_FORMAT_UINT_1X32 = 0x0d, ///< 1 channel, unsigned 32-bit integers
HIP_RES_VIEW_FORMAT_UINT_2X32 = 0x0e, ///< 2 channels, unsigned 32-bit integers
HIP_RES_VIEW_FORMAT_UINT_4X32 = 0x0f, ///< 4 channels, unsigned 32-bit integers
HIP_RES_VIEW_FORMAT_SINT_1X32 = 0x10, ///< 1 channel, signed 32-bit integers
HIP_RES_VIEW_FORMAT_SINT_2X32 = 0x11, ///< 2 channels, signed 32-bit integers
HIP_RES_VIEW_FORMAT_SINT_4X32 = 0x12, ///< 4 channels, signed 32-bit integers
HIP_RES_VIEW_FORMAT_FLOAT_1X16 = 0x13, ///< 1 channel, 16-bit floating point
HIP_RES_VIEW_FORMAT_FLOAT_2X16 = 0x14, ///< 2 channels, 16-bit floating point
HIP_RES_VIEW_FORMAT_FLOAT_4X16 = 0x15, ///< 4 channels, 16-bit floating point
HIP_RES_VIEW_FORMAT_FLOAT_1X32 = 0x16, ///< 1 channel, 32-bit floating point
HIP_RES_VIEW_FORMAT_FLOAT_2X32 = 0x17, ///< 2 channels, 32-bit floating point
HIP_RES_VIEW_FORMAT_FLOAT_4X32 = 0x18, ///< 4 channels, 32-bit floating point
HIP_RES_VIEW_FORMAT_UNSIGNED_BC1 = 0x19, ///< Block-compressed 1
HIP_RES_VIEW_FORMAT_UNSIGNED_BC2 = 0x1a, ///< Block-compressed 2
HIP_RES_VIEW_FORMAT_UNSIGNED_BC3 = 0x1b, ///< Block-compressed 3
HIP_RES_VIEW_FORMAT_UNSIGNED_BC4 = 0x1c, ///< Block-compressed 4 unsigned
HIP_RES_VIEW_FORMAT_SIGNED_BC4 = 0x1d, ///< Block-compressed 4 signed
HIP_RES_VIEW_FORMAT_UNSIGNED_BC5 = 0x1e, ///< Block-compressed 5 unsigned
HIP_RES_VIEW_FORMAT_SIGNED_BC5 = 0x1f, ///< Block-compressed 5 signed
HIP_RES_VIEW_FORMAT_UNSIGNED_BC6H = 0x20, ///< Block-compressed 6 unsigned half-float
HIP_RES_VIEW_FORMAT_SIGNED_BC6H = 0x21, ///< Block-compressed 6 signed half-float
HIP_RES_VIEW_FORMAT_UNSIGNED_BC7 = 0x22 ///< Block-compressed 7
} HIPresourceViewFormat;
/**
* HIP resource descriptor
*/
typedef struct hipResourceDesc {
enum hipResourceType resType; ///< Resource type
union {
struct {
hipArray_t array; ///< HIP array
} array;
struct {
hipMipmappedArray_t mipmap; ///< HIP mipmapped array
} mipmap;
struct {
void* devPtr; ///< Device pointer
struct hipChannelFormatDesc desc; ///< Channel format description
size_t sizeInBytes; ///< Size in bytes
} linear;
struct {
void* devPtr; ///< Device pointer
struct hipChannelFormatDesc desc; ///< Channel format description
size_t width; ///< Width of the array in elements
size_t height; ///< Height of the array in elements
size_t pitchInBytes; ///< Pitch between two rows in bytes
} pitch2D;
} res;
} hipResourceDesc;
/**
* HIP resource view descriptor struct
*/
typedef struct HIP_RESOURCE_DESC_st {
HIPresourcetype resType; ///< Resource type
union {
struct {
hipArray_t hArray; ///< HIP array
} array;
struct {
hipMipmappedArray_t hMipmappedArray; ///< HIP mipmapped array
} mipmap;
struct {
hipDeviceptr_t devPtr; ///< Device pointer
hipArray_Format format; ///< Array format
unsigned int numChannels; ///< Channels per array element
size_t sizeInBytes; ///< Size in bytes
} linear;
struct {
hipDeviceptr_t devPtr; ///< Device pointer
hipArray_Format format; ///< Array format
unsigned int numChannels; ///< Channels per array element
size_t width; ///< Width of the array in elements
size_t height; ///< Height of the array in elements
size_t pitchInBytes; ///< Pitch between two rows in bytes
} pitch2D;
struct {
int reserved[32];
} reserved;
} res;
unsigned int flags; ///< Flags (must be zero)
} HIP_RESOURCE_DESC;
/**
* HIP resource view descriptor
*/
struct hipResourceViewDesc {
enum hipResourceViewFormat format; ///< Resource view format
size_t width; ///< Width of the resource view
size_t height; ///< Height of the resource view
size_t depth; ///< Depth of the resource view
unsigned int firstMipmapLevel; ///< First defined mipmap level
unsigned int lastMipmapLevel; ///< Last defined mipmap level
unsigned int firstLayer; ///< First layer index
unsigned int lastLayer; ///< Last layer index
};
/**
* Resource view descriptor
*/
typedef struct HIP_RESOURCE_VIEW_DESC_st {
HIPresourceViewFormat format; ///< Resource view format
size_t width; ///< Width of the resource view
size_t height; ///< Height of the resource view
size_t depth; ///< Depth of the resource view
unsigned int firstMipmapLevel; ///< First defined mipmap level
unsigned int lastMipmapLevel; ///< Last defined mipmap level
unsigned int firstLayer; ///< First layer index
unsigned int lastLayer; ///< Last layer index
unsigned int reserved[16];
} HIP_RESOURCE_VIEW_DESC;
/**
* Memory copy types
*/
#if !defined(__HIPCC_RTC__)
typedef enum hipMemcpyKind {
hipMemcpyHostToHost = 0, ///< Host-to-Host Copy
hipMemcpyHostToDevice = 1, ///< Host-to-Device Copy
hipMemcpyDeviceToHost = 2, ///< Device-to-Host Copy
hipMemcpyDeviceToDevice = 3, ///< Device-to-Device Copy
hipMemcpyDefault = 4, ///< Runtime will automatically determine
///< copy-kind based on virtual addresses.
hipMemcpyDeviceToDeviceNoCU = 1024 ///< Device-to-Device Copy without using compute units
} hipMemcpyKind;
/**
* HIP pithed pointer
*/
typedef struct hipPitchedPtr {
void* ptr; ///< Pointer to the allocated memory
size_t pitch; ///< Pitch in bytes
size_t xsize; ///< Logical size of the first dimension of allocation in elements
size_t ysize; ///< Logical size of the second dimension of allocation in elements
} hipPitchedPtr;
/**
* HIP extent
*/
typedef struct hipExtent {
size_t width; // Width in elements when referring to array memory, in bytes when referring to
// linear memory
size_t height;
size_t depth;
} hipExtent;
/**
* HIP position
*/
typedef struct hipPos {
size_t x; ///< X coordinate
size_t y; ///< Y coordinate
size_t z; ///< Z coordinate
} hipPos;
/**
* HIP 3D memory copy parameters
*/
typedef struct hipMemcpy3DParms {
hipArray_t srcArray; ///< Source array
struct hipPos srcPos; ///< Source position
struct hipPitchedPtr srcPtr; ///< Source pointer
hipArray_t dstArray; ///< Destination array
struct hipPos dstPos; ///< Destination position
struct hipPitchedPtr dstPtr; ///< Destination pointer
struct hipExtent extent; ///< Extent of 3D memory copy
enum hipMemcpyKind kind; ///< Kind of 3D memory copy
} hipMemcpy3DParms;
/**
* HIP 3D memory copy
*/
typedef struct HIP_MEMCPY3D {
size_t srcXInBytes; ///< Source X in bytes
size_t srcY; ///< Source Y
size_t srcZ; ///< Source Z
size_t srcLOD; ///< Source LOD
hipMemoryType srcMemoryType; ///< Source memory type
const void* srcHost; ///< Source host pointer
hipDeviceptr_t srcDevice; ///< Source device
hipArray_t srcArray; ///< Source array
size_t srcPitch; ///< Source pitch
size_t srcHeight; ///< Source height
size_t dstXInBytes; ///< Destination X in bytes
size_t dstY; ///< Destination Y
size_t dstZ; ///< Destination Z
size_t dstLOD; ///< Destination LOD
hipMemoryType dstMemoryType; ///< Destination memory type
void* dstHost; ///< Destination host pointer
hipDeviceptr_t dstDevice; ///< Destination device
hipArray_t dstArray; ///< Destination array
size_t dstPitch; ///< Destination pitch
size_t dstHeight; ///< Destination height
size_t WidthInBytes; ///< Width in bytes of 3D memory copy
size_t Height; ///< Height in bytes of 3D memory copy
size_t Depth; ///< Depth in bytes of 3D memory copy
} HIP_MEMCPY3D;
/**
* Specifies the type of location
*/
typedef enum hipMemLocationType {
hipMemLocationTypeInvalid = 0,
hipMemLocationTypeNone = 0,
hipMemLocationTypeDevice = 1, ///< Device location, thus it's HIP device ID
hipMemLocationTypeHost = 2, ///< Host location, id is ignored
hipMemLocationTypeHostNuma = 3, ///< Host NUMA node location, id is host NUMA node id
hipMemLocationTypeHostNumaCurrent =
4 ///< Host NUMA node closest to current thread’s CPU, id is ignored
} hipMemLocationType;
/**
* Specifies a memory location.
*
* To specify a gpu, set type = @p hipMemLocationTypeDevice and set id = the gpu's device ID
*/
typedef struct hipMemLocation {
hipMemLocationType type; ///< Specifies the location type, which describes the meaning of id
int id; ///< Identifier for the provided location type @p hipMemLocationType
} hipMemLocation;
/**
* Flags to specify for copies within a batch. Used with hipMemcpyBatchAsync
*/
typedef enum hipMemcpyFlags {
hipMemcpyFlagDefault = 0x0, ///< Default flag
hipMemcpyFlagPreferOverlapWithCompute = 0x1 ///< Tries to overlap copy with compute work.
} hipMemcpyFlags;
/**
* Flags to specify order in which source pointer is accessed by Batch memcpy
*/
typedef enum hipMemcpySrcAccessOrder {
hipMemcpySrcAccessOrderInvalid = 0x0, ///< Default Invalid.
hipMemcpySrcAccessOrderStream = 0x1, ///< Access to source pointer must be in stream order.
hipMemcpySrcAccessOrderDuringApiCall =
0x2, ///< Access to source pointer can be out of stream order and all accesses must be
///< complete before API call returns.
hipMemcpySrcAccessOrderAny =
0x3, ///< Access to the source pointer can be out of stream order and the accesses can happen
///< even after the API call return.
hipMemcpySrcAccessOrderMax = 0x7FFFFFFF
} hipMemcpySrcAccessOrder;
/**
* Attributes for copies within a batch.
*/
typedef struct hipMemcpyAttributes {
hipMemcpySrcAccessOrder
srcAccessOrder; ///< Source access ordering to be observed for copies with this attribute.
hipMemLocation srcLocHint; ///< Location hint for src operand.
hipMemLocation dstLocHint; ///< Location hint for destination operand.
unsigned int flags; ///< Additional Flags for copies. See hipMemcpyFlags.
} hipMemcpyAttributes;
/**
* Operand types for individual copies within a batch
*/
typedef enum hipMemcpy3DOperandType {
hipMemcpyOperandTypePointer = 0x1, ///< Mempcy operand is a valid pointer.
hipMemcpyOperandTypeArray = 0x2, ///< Memcpy operand is a valid hipArray.
hipMemcpyOperandTypeMax = 0x7FFFFFFF
} hipMemcpy3DOperandType;
/**
* Struct representing offset into a hipArray_t in elements.
*/
typedef struct hipOffset3D {
size_t x;
size_t y;
size_t z;
} hipOffset3D;
/**
* Struct representing an operand for copy with hipMemcpy3DBatchAsync.
*/
typedef struct hipMemcpy3DOperand {
hipMemcpy3DOperandType type;
union {
struct {
void* ptr;
size_t rowLength; ///< Length of each row in elements.
size_t layerHeight; ///< Height of each layer in elements.
hipMemLocation locHint; ///< Location Hint for the operand.
} ptr;
struct {
hipArray_t array; ///< Array struct for hipMemcpyOperandTypeArray.
hipOffset3D offset; ///< Offset into array in elements.
} array;
} op;
} hipMemcpy3DOperand;
/**
* HIP 3D Batch Op
*/
typedef struct hipMemcpy3DBatchOp {
hipMemcpy3DOperand src;
hipMemcpy3DOperand dst;
hipExtent extent;
hipMemcpySrcAccessOrder srcAccessOrder;
unsigned int flags;
} hipMemcpy3DBatchOp;
typedef struct hipMemcpy3DPeerParms {
hipArray_t srcArray; ///< Source memory address
hipPos srcPos; ///< Source position offset
hipPitchedPtr srcPtr; ///< Pitched source memory address
int srcDevice; ///< Source device
hipArray_t dstArray; ///< Destination memory address
hipPos dstPos; ///< Destination position offset
hipPitchedPtr dstPtr; ///< Pitched destination memory address
int dstDevice; ///< Destination device
hipExtent extent; ///< Requested memory copy size
} hipMemcpy3DPeerParms;
/**
* @brief Make hipPitchedPtr
*
* @param [in] d Pointer to the allocated memory
* @param [in] p Pitch in bytes
* @param [in] xsz Logical size of the first dimension of allocation in elements
* @param [in] ysz Logical size of the second dimension of allocation in elements
*
* @returns The created hipPitchedPtr
*/
static inline struct hipPitchedPtr make_hipPitchedPtr(void* d, size_t p, size_t xsz, size_t ysz) {
struct hipPitchedPtr s;
s.ptr = d;
s.pitch = p;
s.xsize = xsz;
s.ysize = ysz;
return s;
}
/**
* @brief Make hipPos struct
*
* @param [in] x X coordinate of the new hipPos
* @param [in] y Y coordinate of the new hipPos
* @param [in] z Z coordinate of the new hipPos
*
* @returns The created hipPos struct
*/
static inline struct hipPos make_hipPos(size_t x, size_t y, size_t z) {
struct hipPos p;
p.x = x;
p.y = y;
p.z = z;
return p;
}
/**
* @brief Make hipExtent struct
*
* @param [in] w Width of the new hipExtent
* @param [in] h Height of the new hipExtent
* @param [in] d Depth of the new hipExtent
*
* @returns The created hipExtent struct
*/
static inline struct hipExtent make_hipExtent(size_t w, size_t h, size_t d) {
struct hipExtent e;
e.width = w;
e.height = h;
e.depth = d;
return e;
}
typedef enum hipFunction_attribute {
HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, ///< The maximum number of threads per block. Depends
///< on function and device.
HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES, ///< The statically allocated shared memory size in bytes
///< per block required by the function.
HIP_FUNC_ATTRIBUTE_CONST_SIZE_BYTES, ///< The user-allocated constant memory by the function in
///< bytes.
HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES, ///< The local memory usage of each thread by this function
///< in bytes.
HIP_FUNC_ATTRIBUTE_NUM_REGS, ///< The number of registers used by each thread of this function.
HIP_FUNC_ATTRIBUTE_PTX_VERSION, ///< PTX version
HIP_FUNC_ATTRIBUTE_BINARY_VERSION, ///< Binary version
HIP_FUNC_ATTRIBUTE_CACHE_MODE_CA, ///< Cache mode
HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, ///< The maximum dynamic shared memory per
///< block for this function in bytes.
HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT, ///< The shared memory carveout preference
///< in percent of the maximum shared
///< memory.
HIP_FUNC_ATTRIBUTE_MAX
} hipFunction_attribute;
typedef enum hipPointer_attribute {
HIP_POINTER_ATTRIBUTE_CONTEXT = 1, ///< The context on which a pointer was allocated
///< @warning This attribute is not supported in HIP
HIP_POINTER_ATTRIBUTE_MEMORY_TYPE, ///< memory type describing the location of a pointer
HIP_POINTER_ATTRIBUTE_DEVICE_POINTER, ///< address at which the pointer is allocated on the
///< device
HIP_POINTER_ATTRIBUTE_HOST_POINTER, ///< address at which the pointer is allocated on the host
HIP_POINTER_ATTRIBUTE_P2P_TOKENS, ///< A pair of tokens for use with Linux kernel interface
///< @warning This attribute is not supported in HIP
HIP_POINTER_ATTRIBUTE_SYNC_MEMOPS, ///< Synchronize every synchronous memory operation
///< initiated on this region
HIP_POINTER_ATTRIBUTE_BUFFER_ID, ///< Unique ID for an allocated memory region
HIP_POINTER_ATTRIBUTE_IS_MANAGED, ///< Indicates if the pointer points to managed memory
HIP_POINTER_ATTRIBUTE_DEVICE_ORDINAL, ///< device ordinal of a device on which a pointer
///< was allocated or registered
HIP_POINTER_ATTRIBUTE_IS_LEGACY_HIP_IPC_CAPABLE, ///< if this pointer maps to an allocation
///< that is suitable for hipIpcGetMemHandle
///< @warning This attribute is not supported in
///< HIP
HIP_POINTER_ATTRIBUTE_RANGE_START_ADDR, ///< Starting address for this requested pointer
HIP_POINTER_ATTRIBUTE_RANGE_SIZE, ///< Size of the address range for this requested pointer
HIP_POINTER_ATTRIBUTE_MAPPED, ///< tells if this pointer is in a valid address range
///< that is mapped to a backing allocation
HIP_POINTER_ATTRIBUTE_ALLOWED_HANDLE_TYPES, ///< Bitmask of allowed hipmemAllocationHandleType
///< for this allocation @warning This attribute is
///< not supported in HIP
HIP_POINTER_ATTRIBUTE_IS_GPU_DIRECT_RDMA_CAPABLE, ///< returns if the memory referenced by
///< this pointer can be used with the
///< GPUDirect RDMA API
///< @warning This attribute is not supported
///< in HIP
HIP_POINTER_ATTRIBUTE_ACCESS_FLAGS, ///< Returns the access flags the device associated with
///< for the corresponding memory referenced by the ptr
HIP_POINTER_ATTRIBUTE_MEMPOOL_HANDLE ///< Returns the mempool handle for the allocation if
///< it was allocated from a mempool
///< @warning This attribute is not supported in HIP
} hipPointer_attribute;
// doxygen end DriverTypes
/**
* @}
*/
#endif // !defined(__HIPCC_RTC__)
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif

36
external/hip/hip/hip_bf16.h vendored Normal file
View File

@@ -0,0 +1,36 @@
/*
Copyright (c) 2023 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 HIP_INCLUDE_HIP_HIP_BF16_H
#define HIP_INCLUDE_HIP_HIP_BF16_H
#include <hip/hip_common.h>
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/amd_detail/amd_hip_bf16.h>
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/nvidia_detail/nvidia_hip_bf16.h>
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif // HIP_INCLUDE_HIP_HIP_BF16_H

44
external/hip/hip/hip_bfloat16.h vendored Normal file
View File

@@ -0,0 +1,44 @@
/**
* MIT License
*
* Copyright (c) 2019 - 2022 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.
*/
/*!\file
* \brief hip_bfloat16.h provides struct for hip_bfloat16 typedef
*/
#ifndef _HIP_BFLOAT16_H_
#define _HIP_BFLOAT16_H_
#if !defined(__HIPCC_RTC__)
#include <hip/hip_common.h>
#endif
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/amd_detail/amd_hip_bfloat16.h>
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#warning "hip_bfloat16.h is not supported on nvidia platform"
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif // _HIP_BFLOAT16_H_

100
external/hip/hip/hip_common.h vendored Normal file
View File

@@ -0,0 +1,100 @@
/*
Copyright (c) 2015 - 2023 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 HIP_INCLUDE_HIP_HIP_COMMON_H
#define HIP_INCLUDE_HIP_HIP_COMMON_H
#if defined(__clang__)
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wreserved-macro-identifier"
#endif
// Common code included at start of every hip file.
// Auto enable __HIP_PLATFORM_AMD__ if compiling on AMD platform
// Other compiler (GCC,ICC,etc) need to set one of these macros explicitly
#if defined(__clang__) && defined(__HIP__)
#ifndef __HIP_PLATFORM_AMD__
#define __HIP_PLATFORM_AMD__
#endif
#endif // defined(__clang__) && defined(__HIP__)
// Auto enable __HIP_PLATFORM_NVIDIA__ if compiling with NVIDIA platform
#if defined(__NVCC__) || (defined(__clang__) && defined(__CUDA__) && !defined(__HIP__))
#ifndef __HIP_PLATFORM_NVIDIA__
#define __HIP_PLATFORM_NVIDIA__
#endif
#ifdef __CUDACC__
#define __HIPCC__
#endif
#endif //__NVCC__
// Auto enable __HIP_DEVICE_COMPILE__ if compiled in HCC or NVCC device path
#if (defined(__HCC_ACCELERATOR__) && __HCC_ACCELERATOR__ != 0) || \
(defined(__CUDA_ARCH__) && __CUDA_ARCH__ != 0)
#define __HIP_DEVICE_COMPILE__ 1
#endif
#ifdef __GNUC__
#define HIP_PUBLIC_API __attribute__((visibility("default")))
#define HIP_INTERNAL_EXPORTED_API __attribute__((visibility("default")))
#else
#define HIP_PUBLIC_API
#define HIP_INTERNAL_EXPORTED_API
#endif
#if __HIP_DEVICE_COMPILE__ == 0
// 32-bit Atomics
#define __HIP_ARCH_HAS_GLOBAL_INT32_ATOMICS__ (0)
#define __HIP_ARCH_HAS_GLOBAL_FLOAT_ATOMIC_EXCH__ (0)
#define __HIP_ARCH_HAS_SHARED_INT32_ATOMICS__ (0)
#define __HIP_ARCH_HAS_SHARED_FLOAT_ATOMIC_EXCH__ (0)
#define __HIP_ARCH_HAS_FLOAT_ATOMIC_ADD__ (0)
// 64-bit Atomics
#define __HIP_ARCH_HAS_GLOBAL_INT64_ATOMICS__ (0)
#define __HIP_ARCH_HAS_SHARED_INT64_ATOMICS__ (0)
// Doubles
#define __HIP_ARCH_HAS_DOUBLES__ (0)
// Warp cross-lane operations
#define __HIP_ARCH_HAS_WARP_VOTE__ (0)
#define __HIP_ARCH_HAS_WARP_BALLOT__ (0)
#define __HIP_ARCH_HAS_WARP_SHUFFLE__ (0)
#define __HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__ (0)
// Sync
#define __HIP_ARCH_HAS_THREAD_FENCE_SYSTEM__ (0)
#define __HIP_ARCH_HAS_SYNC_THREAD_EXT__ (0)
// Misc
#define __HIP_ARCH_HAS_SURFACE_FUNCS__ (0)
#define __HIP_ARCH_HAS_3DGRID__ (0)
#define __HIP_ARCH_HAS_DYNAMIC_PARALLEL__ (0)
#endif
#if defined(__clang__)
#pragma clang diagnostic pop
#endif
#endif

38
external/hip/hip/hip_complex.h vendored Normal file
View File

@@ -0,0 +1,38 @@
/*
Copyright (c) 2015 - 2023 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 HIP_INCLUDE_HIP_HIP_COMPLEX_H
#define HIP_INCLUDE_HIP_HIP_COMPLEX_H
#if !defined(__HIPCC_RTC__)
#include <hip/hip_common.h>
#endif
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/amd_detail/amd_hip_complex.h>
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/nvidia_detail/nvidia_hip_complex.h>
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif

46
external/hip/hip/hip_cooperative_groups.h vendored Normal file
View File

@@ -0,0 +1,46 @@
/*
Copyright (c) 2015 - 2021 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.
*/
/**
* @file hip_cooperative_groups.h
*
* @brief Defines new types and device API wrappers for `Cooperative Group`
* feature.
*/
#ifndef HIP_INCLUDE_HIP_HIP_COOPERATIVE_GROUP_H
#define HIP_INCLUDE_HIP_HIP_COOPERATIVE_GROUP_H
#include <hip/hip_version.h>
#include <hip/hip_common.h>
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#if __cplusplus && defined(__clang__) && defined(__HIP__)
#include <hip/amd_detail/amd_hip_cooperative_groups.h>
#endif
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/nvidia_detail/nvidia_hip_cooperative_groups.h>
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif // HIP_INCLUDE_HIP_HIP_COOPERATIVE_GROUP_H

119
external/hip/hip/hip_deprecated.h vendored Normal file
View File

@@ -0,0 +1,119 @@
/*
* Copyright (C) Advanced Micro Devices, Inc.
*
* 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 COPYRIGHT HOLDER(S) 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.
*/
#pragma once
// This file will add older hip functions used in the versioning system
// Find the deprecated functions and structs in hip_device.cpp
// This struct is also kept in hip_device.cpp
typedef struct hipDeviceProp_tR0000 {
char name[256]; ///< Device name.
size_t totalGlobalMem; ///< Size of global memory region (in bytes).
size_t sharedMemPerBlock; ///< Size of shared memory region (in bytes).
int regsPerBlock; ///< Registers per block.
int warpSize; ///< Warp size.
int maxThreadsPerBlock; ///< Max work items per work group or workgroup max size.
int maxThreadsDim[3]; ///< Max number of threads in each dimension (XYZ) of a block.
int maxGridSize[3]; ///< Max grid dimensions (XYZ).
int clockRate; ///< Max clock frequency of the multiProcessors in khz.
int memoryClockRate; ///< Max global memory clock frequency in khz.
int memoryBusWidth; ///< Global memory bus width in bits.
size_t totalConstMem; ///< Size of shared memory region (in bytes).
int major; ///< Major compute capability. On HCC, this is an approximation and features may
///< differ from CUDA CC. See the arch feature flags for portable ways to query
///< feature caps.
int minor; ///< Minor compute capability. On HCC, this is an approximation and features may
///< differ from CUDA CC. See the arch feature flags for portable ways to query
///< feature caps.
int multiProcessorCount; ///< Number of multi-processors. When the GPU works in Compute
///< Unit (CU) mode, this value equals the number of CUs;
///< when in Workgroup Processor (WGP) mode, this value equels
///< half of CUs, because a single WGP contains two CUs.
int l2CacheSize; ///< L2 cache size.
int maxThreadsPerMultiProcessor; ///< Maximum resident threads per multi-processor.
int computeMode; ///< Compute mode.
int clockInstructionRate; ///< Frequency in khz of the timer used by the device-side "clock*"
///< instructions. New for HIP.
hipDeviceArch_t arch; ///< Architectural feature flags. New for HIP.
int concurrentKernels; ///< Device can possibly execute multiple kernels concurrently.
int pciDomainID; ///< PCI Domain ID
int pciBusID; ///< PCI Bus ID.
int pciDeviceID; ///< PCI Device ID.
size_t maxSharedMemoryPerMultiProcessor; ///< Maximum Shared Memory Per Multiprocessor.
int isMultiGpuBoard; ///< 1 if device is on a multi-GPU board, 0 if not.
int canMapHostMemory; ///< Check whether HIP can map host memory
int gcnArch; ///< DEPRECATED: use gcnArchName instead
char gcnArchName[256]; ///< AMD GCN Arch Name.
int integrated; ///< APU vs dGPU
int cooperativeLaunch; ///< HIP device supports cooperative launch
int cooperativeMultiDeviceLaunch; ///< HIP device supports cooperative launch on multiple
///< devices
int maxTexture1DLinear; ///< Maximum size for 1D textures bound to linear memory
int maxTexture1D; ///< Maximum number of elements in 1D images
int maxTexture2D[2]; ///< Maximum dimensions (width, height) of 2D images, in image elements
int maxTexture3D[3]; ///< Maximum dimensions (width, height, depth) of 3D images, in image
///< elements
unsigned int* hdpMemFlushCntl; ///< Addres of HDP_MEM_COHERENCY_FLUSH_CNTL register
unsigned int* hdpRegFlushCntl; ///< Addres of HDP_REG_COHERENCY_FLUSH_CNTL register
size_t memPitch; ///< Maximum pitch in bytes allowed by memory copies
size_t textureAlignment; ///< Alignment requirement for textures
size_t texturePitchAlignment; ///< Pitch alignment requirement for texture references bound to
///< pitched memory
int kernelExecTimeoutEnabled; ///< Run time limit for kernels executed on the device
int ECCEnabled; ///< Device has ECC support enabled
int tccDriver; ///< 1:If device is Tesla device using TCC driver, else 0
int cooperativeMultiDeviceUnmatchedFunc; ///< HIP device supports cooperative launch on
///< multiple
/// devices with unmatched functions
int cooperativeMultiDeviceUnmatchedGridDim; ///< HIP device supports cooperative launch on
///< multiple
/// devices with unmatched grid dimensions
int cooperativeMultiDeviceUnmatchedBlockDim; ///< HIP device supports cooperative launch on
///< multiple
/// devices with unmatched block dimensions
int cooperativeMultiDeviceUnmatchedSharedMem; ///< HIP device supports cooperative launch on
///< multiple
/// devices with unmatched shared memories
int isLargeBar; ///< 1: if it is a large PCI bar device, else 0
int asicRevision; ///< Revision of the GPU in this device
int managedMemory; ///< Device supports allocating managed memory on this system
int directManagedMemAccessFromHost; ///< Host can directly access managed memory on the device
///< without migration
int concurrentManagedAccess; ///< Device can coherently access managed memory concurrently with
///< the CPU
int pageableMemoryAccess; ///< Device supports coherently accessing pageable memory
///< without calling hipHostRegister on it
int pageableMemoryAccessUsesHostPageTables; ///< Device accesses pageable memory via the host's
///< page tables
} hipDeviceProp_tR0000;
#ifdef __cplusplus
extern "C" {
#endif
hipError_t hipGetDevicePropertiesR0000(hipDeviceProp_tR0000* prop, int device);
hipError_t hipChooseDeviceR0000(int* device, const hipDeviceProp_tR0000* prop);
#ifdef __cplusplus
}
#endif

162
external/hip/hip/hip_ext.h vendored Normal file
View File

@@ -0,0 +1,162 @@
/*
Copyright (c) 2015 - 2021 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 HIP_INCLUDE_HIP_HIP_EXT_H
#define HIP_INCLUDE_HIP_HIP_EXT_H
#include "hip/hip_runtime.h"
#if defined(__cplusplus)
#include <tuple>
#include <type_traits>
#endif
/** @addtogroup Execution Execution Control
* @{
*/
/**
* @brief Launches kernel with parameters and shared memory on stream with arguments passed
* to kernel params or extra arguments.
*
* @param [in] f Kernel to launch.
* @param [in] globalWorkSizeX X grid dimension specified in work-items.
* @param [in] globalWorkSizeY Y grid dimension specified in work-items.
* @param [in] globalWorkSizeZ Z grid dimension specified in work-items.
* @param [in] localWorkSizeX X block dimension specified in work-items.
* @param [in] localWorkSizeY Y block dimension specified in work-items.
* @param [in] localWorkSizeZ Z block dimension specified in work-items.
* @param [in] sharedMemBytes Amount of dynamic shared memory to allocate for this kernel.
* HIP-Clang compiler provides support for extern shared declarations.
* @param [in] hStream Stream where the kernel should be dispatched.
* May be 0, in which case the default stream is used with associated synchronization rules.
* @param [in] kernelParams pointer to kernel parameters.
* @param [in] extra Pointer to kernel arguments. These are passed directly to the kernel and
* must be in the memory layout and alignment expected by the kernel.
* All passed arguments must be naturally aligned according to their type. The memory address of
* each argument should be a multiple of its size in bytes. Please refer to
* hip_porting_driver_api.md for sample usage.
* @param [in] startEvent If non-null, specified event will be updated to track the start time of
* the kernel launch. The event must be created before calling this API.
* @param [in] stopEvent If non-null, specified event will be updated to track the stop time of
* the kernel launch. The event must be created before calling this API.
* @param [in] flags The value of hipExtAnyOrderLaunch, signifies if kernel can be
* launched in any order.
* @returns #hipSuccess, #hipInvalidDeviceId, #hipErrorNotInitialized, #hipErrorInvalidValue.
*
* HIP/ROCm actually updates the start event when the associated kernel completes.
* Currently, timing between startEvent and stopEvent does not include the time it takes to perform
* a system scope release/cache flush - only the time it takes to issues writes to cache.
*
* @note For this HIP API, the flag 'hipExtAnyOrderLaunch' is not supported on AMD GFX9xx boards.
*
*/
HIP_PUBLIC_API
extern "C" hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
uint32_t localWorkSizeX, uint32_t localWorkSizeY,
uint32_t localWorkSizeZ, size_t sharedMemBytes,
hipStream_t hStream, void** kernelParams,
void** extra, hipEvent_t startEvent __dparm(NULL),
hipEvent_t stopEvent __dparm(NULL),
uint32_t flags __dparm(0));
/**
* @brief This HIP API is deprecated, please use hipExtModuleLaunchKernel() instead.
*
*/
HIP_DEPRECATED("use hipExtModuleLaunchKernel instead")
HIP_PUBLIC_API
extern "C" hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
uint32_t localWorkSizeX, uint32_t localWorkSizeY,
uint32_t localWorkSizeZ, size_t sharedMemBytes,
hipStream_t hStream, void** kernelParams,
void** extra, hipEvent_t startEvent __dparm(NULL),
hipEvent_t stopEvent __dparm(NULL));
#if defined(__cplusplus)
/**
* @brief Launches kernel from the pointer address, with arguments and shared memory on stream.
*
* @param [in] function_address pointer to the Kernel to launch.
* @param [in] numBlocks number of blocks.
* @param [in] dimBlocks dimension of a block.
* @param [in] args pointer to kernel arguments.
* @param [in] sharedMemBytes Amount of dynamic shared memory to allocate for this kernel.
* HIP-Clang compiler provides support for extern shared declarations.
* @param [in] stream Stream where the kernel should be dispatched.
* May be 0, in which case the default stream is used with associated synchronization rules.
* @param [in] startEvent If non-null, specified event will be updated to track the start time of
* the kernel launch. The event must be created before calling this API.
* @param [in] stopEvent If non-null, specified event will be updated to track the stop time of
* the kernel launch. The event must be created before calling this API.
* @param [in] flags The value of hipExtAnyOrderLaunch, signifies if kernel can be
* launched in any order.
* @returns #hipSuccess, #hipInvalidDeviceId, #hipErrorNotInitialized, #hipErrorInvalidValue.
*
*/
extern "C" hipError_t hipExtLaunchKernel(const void* function_address, dim3 numBlocks,
dim3 dimBlocks, void** args, size_t sharedMemBytes,
hipStream_t stream, hipEvent_t startEvent,
hipEvent_t stopEvent, int flags);
/**
* @brief Launches kernel with dimention parameters and shared memory on stream with templated
* kernel and arguments.
*
* @param [in] kernel Kernel to launch.
* @param [in] numBlocks const number of blocks.
* @param [in] dimBlocks const dimension of a block.
* @param [in] sharedMemBytes Amount of dynamic shared memory to allocate for this kernel.
* HIP-Clang compiler provides support for extern shared declarations.
* @param [in] stream Stream where the kernel should be dispatched.
* May be 0, in which case the default stream is used with associated synchronization rules.
* @param [in] startEvent If non-null, specified event will be updated to track the start time of
* the kernel launch. The event must be created before calling this API.
* @param [in] stopEvent If non-null, specified event will be updated to track the stop time of
* the kernel launch. The event must be created before calling this API.
* @param [in] flags The value of hipExtAnyOrderLaunch, signifies if kernel can be
* launched in any order.
* @param [in] args templated kernel arguments.
*
*/
template <typename... Args, typename F = void (*)(Args...)>
inline void hipExtLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
std::uint32_t sharedMemBytes, hipStream_t stream,
hipEvent_t startEvent, hipEvent_t stopEvent, std::uint32_t flags,
Args... args) {
constexpr size_t count = sizeof...(Args);
auto tup_ = std::tuple<Args...>{args...};
auto tup = validateArgsCountType(kernel, tup_);
void* _Args[count];
pArgs<0>(tup, _Args);
auto k = reinterpret_cast<void*>(kernel);
hipExtLaunchKernel(k, numBlocks, dimBlocks, _Args, sharedMemBytes, stream, startEvent, stopEvent,
(int)flags);
}
#endif // defined(__cplusplus)
// doxygen end AMD-specific features
/**
* @}
*/
#endif // #iidef HIP_INCLUDE_HIP_HIP_EXT_H

31
external/hip/hip/hip_ext_ocp.h vendored Normal file
View File

@@ -0,0 +1,31 @@
/*
Copyright (c) 2024 - present 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 HIP_INCLUDE_HIP_HIP_EXT_OCP_H
#define HIP_INCLUDE_HIP_HIP_EXT_OCP_H
#include <hip/amd_detail/amd_hip_ocp_types.h>
#include <hip/amd_detail/amd_hip_ocp_fp.hpp>
#include <hip/amd_detail/amd_hip_ocp_fp_cxx.hpp>
#endif // HIP_INCLUDE_HIP_HIP_EXT_OCP_H

44
external/hip/hip/hip_fp16.h vendored Normal file
View File

@@ -0,0 +1,44 @@
/*
Copyright (c) 2015 - 2021 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 HIP_INCLUDE_HIP_HIP_FP16_H
#define HIP_INCLUDE_HIP_HIP_FP16_H
#include <hip/hip_common.h>
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/amd_detail/amd_hip_fp16.h>
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#define HIPRT_INF_FP16 CUDART_INF_FP16
#define HIPRT_MAX_NORMAL_FP16 CUDART_MAX_NORMAL_FP16
#define HIPRT_MIN_DENORM_FP16 CUDART_MIN_DENORM_FP16
#define HIPRT_NAN_FP16 CUDART_NAN_FP16
#define HIPRT_NEG_ZERO_FP16 CUDART_NEG_ZERO_FP16
#define HIPRT_ONE_FP16 CUDART_ONE_FP16
#define HIPRT_ZERO_FP16 CUDART_ZERO_FP16
#include "cuda_fp16.h"
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif

32
external/hip/hip/hip_fp4.h vendored Normal file
View File

@@ -0,0 +1,32 @@
/*
Copyright (c) 2025 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 HIP_INCLUDE_HIP_HIP_FP4_H
#define HIP_INCLUDE_HIP_HIP_FP4_H
#include <hip/hip_common.h>
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/amd_detail/amd_hip_fp4.h>
#endif
#endif // HIP_INCLUDE_HIP_HIP_FP4_H

32
external/hip/hip/hip_fp6.h vendored Normal file
View File

@@ -0,0 +1,32 @@
/*
Copyright (c) 2025 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 HIP_INCLUDE_HIP_HIP_FP6_H
#define HIP_INCLUDE_HIP_HIP_FP6_H
#include <hip/hip_common.h>
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/amd_detail/amd_hip_fp6.h>
#endif
#endif // HIP_INCLUDE_HIP_HIP_FP6_H

33
external/hip/hip/hip_fp8.h vendored Normal file
View File

@@ -0,0 +1,33 @@
/*
Copyright (c) 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 HIP_INCLUDE_HIP_HIP_FP8_H
#define HIP_INCLUDE_HIP_HIP_FP8_H
#include <hip/hip_common.h>
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
// We only have fnuz defs for now, which are not supported by other platforms
#include <hip/amd_detail/amd_hip_fp8.h>
#endif
#endif // HIP_INCLUDE_HIP_HIP_FP8_H

32
external/hip/hip/hip_gl_interop.h vendored Normal file
View File

@@ -0,0 +1,32 @@
/*
Copyright (c) 2023 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 HIP_GL_INTEROP_H
#define HIP_GL_INTEROP_H
#include <hip/hip_common.h>
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#include "hip/amd_detail/amd_hip_gl_interop.h"
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include "hip/nvidia_detail/nvidia_hip_gl_interop.h"
#endif
#endif

24
external/hip/hip/hip_hcc.h vendored Normal file
View File

@@ -0,0 +1,24 @@
/*
Copyright (c) 2015 - 2021 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 HIP_INCLUDE_HIP_HIP_HCC_H
#define HIP_INCLUDE_HIP_HIP_HCC_H
#warning "hip/hip_hcc.h is deprecated, please use hip/hip_ext.h"
#include "hip/hip_ext.h"
#endif // #ifdef HIP_INCLUDE_HIP_HIP_HCC_H

36
external/hip/hip/hip_math_constants.h vendored Normal file
View File

@@ -0,0 +1,36 @@
/*
Copyright (c) 2015 - 2022 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 HIP_MATH_CONSTANTS_H
#define HIP_MATH_CONSTANTS_H
#if !defined(__HIPCC_RTC__)
#include <hip/hip_common.h>
#endif
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#include "hip/amd_detail/amd_hip_math_constants.h"
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include "hip/nvidia_detail/nvidia_hip_math_constants.h"
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif

27
external/hip/hip/hip_profile.h vendored Normal file
View File

@@ -0,0 +1,27 @@
/*
Copyright (c) 2015 - 2021 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 HIP_INCLUDE_HIP_HIP_PROFILE_H
#define HIP_INCLUDE_HIP_HIP_PROFILE_H
#define HIP_SCOPED_MARKER(markerName, group)
#define HIP_BEGIN_MARKER(markerName, group)
#define HIP_END_MARKER()
#endif

70
external/hip/hip/hip_runtime.h vendored Normal file
View File

@@ -0,0 +1,70 @@
/*
Copyright (c) 2015 - 2025 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.
*/
//! HIP = Heterogeneous-compute Interface for Portability
//!
//! Define a extremely thin runtime layer that allows source code to be compiled unmodified
//! through either AMD CLANG or NVCC. Key features tend to be in the spirit
//! and terminology of CUDA, but with a portable path to other accelerators as well:
//
//! Both paths support rich C++ features including classes, templates, lambdas, etc.
//! Runtime API is C
//! Memory management is based on pure pointers and resembles malloc/free/copy.
//
//! hip_runtime.h : includes everything in hip_api.h, plus math builtins and kernel launch
//! macros. hip_runtime_api.h : Defines HIP API. This is a C header file and does not use any C++
//! features.
#ifndef HIP_INCLUDE_HIP_HIP_RUNTIME_H
#define HIP_INCLUDE_HIP_HIP_RUNTIME_H
#if !defined(__HIPCC_RTC__)
// Some standard header files, these are included by hc.hpp and so want to make them avail on both
// paths to provide a consistent include env and avoid "missing symbol" errors that only appears
// on NVCC path:
#if __cplusplus
#include <cstdint>
#include <cstdlib>
#else
#include <stdint.h>
#include <stdlib.h>
#endif // __cplusplus
#endif // !defined(__HIPCC_RTC__)
#include <hip/hip_version.h>
#include <hip/hip_common.h>
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/amd_detail/amd_hip_runtime.h>
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/nvidia_detail/nvidia_hip_runtime.h>
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#if !defined(__HIPCC_RTC__)
#include <hip/hip_runtime_api.h>
#include <hip/library_types.h>
#endif // !defined(__HIPCC_RTC__)
#include <hip/hip_vector_types.h>
#endif

10431
external/hip/hip/hip_runtime_api.h vendored Normal file
View File

File diff suppressed because it is too large Load Diff

29
external/hip/hip/hip_texture_types.h vendored Normal file
View File

@@ -0,0 +1,29 @@
/*
Copyright (c) 2015 - 2021 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 HIP_INCLUDE_HIP_HIP_TEXTURE_TYPES_H
#define HIP_INCLUDE_HIP_HIP_TEXTURE_TYPES_H
#include <hip/texture_types.h>
#endif

41
external/hip/hip/hip_vector_types.h vendored Normal file
View File

@@ -0,0 +1,41 @@
/*
Copyright (c) 2015 - 2021 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.
*/
//! hip_vector_types.h : Defines the HIP vector types.
#ifndef HIP_INCLUDE_HIP_HIP_VECTOR_TYPES_H
#define HIP_INCLUDE_HIP_HIP_VECTOR_TYPES_H
#include <hip/hip_common.h>
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#if __cplusplus
#include <hip/amd_detail/amd_hip_vector_types.h>
#endif
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include <vector_types.h>
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif

473
external/hip/hip/hiprtc.h vendored Normal file
View File

@@ -0,0 +1,473 @@
/*
Copyright (c) 2015 - 2023 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.
*/
#pragma once
#include <hip/hip_common.h>
#include <hip/linker_types.h>
#if !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/nvidia_detail/nvidia_hiprtc.h>
#elif defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#ifdef __cplusplus
#include <cstdlib>
#else
#include <stdlib.h>
#endif
#ifdef __cplusplus
extern "C" {
#endif /* __cplusplus */
#if !defined(_WIN32)
#pragma GCC visibility push(default)
#endif
/**
*
* @addtogroup GlobalDefs
* @{
*
*/
/**
* hiprtc error code
*/
typedef enum hiprtcResult {
HIPRTC_SUCCESS = 0, ///< Success
HIPRTC_ERROR_OUT_OF_MEMORY = 1, ///< Out of memory
HIPRTC_ERROR_PROGRAM_CREATION_FAILURE = 2, ///< Failed to create program
HIPRTC_ERROR_INVALID_INPUT = 3, ///< Invalid input
HIPRTC_ERROR_INVALID_PROGRAM = 4, ///< Invalid program
HIPRTC_ERROR_INVALID_OPTION = 5, ///< Invalid option
HIPRTC_ERROR_COMPILATION = 6, ///< Compilation error
HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE = 7, ///< Failed in builtin operation
HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = 8, ///< No name expression after compilation
HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = 9, ///< No lowered names before compilation
HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID = 10, ///< Invalid name expression
HIPRTC_ERROR_INTERNAL_ERROR = 11, ///< Internal error
HIPRTC_ERROR_LINKING = 100 ///< Error in linking
} hiprtcResult;
/**
* hiprtc JIT option
*/
#define hiprtcJIT_option hipJitOption
#define HIPRTC_JIT_MAX_REGISTERS \
hipJitOptionMaxRegisters ///< CUDA Only Maximum registers may be used in a
///< thread, passed to compiler
#define HIPRTC_JIT_THREADS_PER_BLOCK \
hipJitOptionThreadsPerBlock ///< CUDA Only Number of thread per block
#define HIPRTC_JIT_WALL_TIME hipJitOptionWallTime ///< CUDA Only Value for total wall clock time
#define HIPRTC_JIT_INFO_LOG_BUFFER \
hipJitOptionInfoLogBuffer ///< CUDA Only Pointer to the buffer with
///< logged information
#define HIPRTC_JIT_INFO_LOG_BUFFER_SIZE_BYTES \
hipJitOptionInfoLogBufferSizeBytes ///< CUDA Only Size of the buffer
///< in bytes for logged info
#define HIPRTC_JIT_ERROR_LOG_BUFFER \
hipJitOptionErrorLogBuffer ///< CUDA Only Pointer to the buffer
///< with logged error(s)
#define HIPRTC_JIT_ERROR_LOG_BUFFER_SIZE_BYTES \
hipJitOptionErrorLogBufferSizeBytes ///< CUDA Only Size of the buffer in
///< bytes for logged error(s)
#define HIPRTC_JIT_OPTIMIZATION_LEVEL \
hipJitOptionOptimizationLevel ///< Value of optimization level for
///< generated codes, acceptable
///< options -O0, -O1, -O2, -O3
#define HIPRTC_JIT_TARGET_FROM_HIPCONTEXT \
hipJitOptionTargetFromContext ///< CUDA Only The target context,
///< which is the default
#define HIPRTC_JIT_TARGET hipJitOptionTarget ///< CUDA Only JIT target
#define HIPRTC_JIT_FALLBACK_STRATEGY hipJitOptionFallbackStrategy ///< CUDA Only Fallback strategy
#define HIPRTC_JIT_GENERATE_DEBUG_INFO \
hipJitOptionGenerateDebugInfo ///< CUDA Only Generate debug information
#define HIPRTC_JIT_LOG_VERBOSE hipJitOptionLogVerbose ///< CUDA Only Generate log verbose
#define HIPRTC_JIT_GENERATE_LINE_INFO \
hipJitOptionGenerateLineInfo ///< CUDA Only Generate line number information
#define HIPRTC_JIT_CACHE_MODE hipJitOptionCacheMode ///< CUDA Only Set cache mode
#define HIPRTC_JIT_NEW_SM3X_OPT hipJitOptionSm3xOpt ///< @deprecated CUDA Only New SM3X option.
#define HIPRTC_JIT_FAST_COMPILE hipJitOptionFastCompile ///< CUDA Only Set fast compile
#define HIPRTC_JIT_GLOBAL_SYMBOL_NAMES \
hipJitOptionGlobalSymbolNames ///< CUDA Only Array of device symbol names to be
///< relocated to the host
#define HIPRTC_JIT_GLOBAL_SYMBOL_ADDRESS \
hipJitOptionGlobalSymbolAddresses ///< CUDA Only Array of host addresses to be
///< relocated to the device
#define HIPRTC_JIT_GLOBAL_SYMBOL_COUNT \
hipJitOptionGlobalSymbolCount ///< CUDA Only Number of symbol count.
#define HIPRTC_JIT_LTO \
hipJitOptionLto ///< @deprecated CUDA Only Enable link-time
///< optimization for device code
#define HIPRTC_JIT_FTZ \
hipJitOptionFtz ///< @deprecated CUDA Only Set
///< single-precision denormals.
#define HIPRTC_JIT_PREC_DIV \
hipJitOptionPrecDiv ///< @deprecated CUDA Only Set
///< single-precision floating-point division
///< and reciprocals
#define HIPRTC_JIT_PREC_SQRT \
hipJitOptionPrecSqrt ///< @deprecated CUDA Only Set
///< single-precision floating-point
///< square root
#define HIPRTC_JIT_FMA \
hipJitOptionFma ///< @deprecated CUDA Only Enable
///< floating-point multiplies and
///< adds/subtracts operations
#define HIPRTC_JIT_POSITION_INDEPENDENT_CODE \
hipJitOptionPositionIndependentCode ///< CUDA Only Generates
///< Position Independent code
#define HIPRTC_JIT_MIN_CTA_PER_SM \
hipJitOptionMinCTAPerSM ///< CUDA Only Hints to JIT compiler
///< the minimum number of CTAs frin
///< kernel's grid to be mapped to SM
#define HIPRTC_JIT_MAX_THREADS_PER_BLOCK \
hipJitOptionMaxThreadsPerBlock ///< CUDA only Maximum number of
///< threads in a thread block
#define HIPRTC_JIT_OVERRIDE_DIRECT_VALUES \
hipJitOptionOverrideDirectiveValues ///< CUDA only Override Directive
///< Values
#define HIPRTC_JIT_NUM_OPTIONS hipJitOptionNumOptions ///< Number of options
#define HIPRTC_JIT_IR_TO_ISA_OPT_EXT \
hipJitOptionIRtoISAOptExt ///< HIP Only Linker options to be
///< passed on to compiler
#define HIPRTC_JIT_IR_TO_ISA_OPT_COUNT_EXT \
hipJitOptionIRtoISAOptCountExt ///< HIP Only Count of linker options
///< to be passed on to
/**
* hiprtc JIT input type
*/
#define hiprtcJITInputType hipJitInputType
#define HIPRTC_JIT_INPUT_CUBIN hipJitInputCubin ///< Cuda only Input Cubin
#define HIPRTC_JIT_INPUT_PTX hipJitInputPtx ///< Cuda only Input PTX
#define HIPRTC_JIT_INPUT_FATBINARY hipJitInputFatBinary ///< Cuda Only Input FAT Binary
#define HIPRTC_JIT_INPUT_OBJECT \
hipJitInputObject ///< Cuda Only Host Object with embedded device code
#define HIPRTC_JIT_INPUT_LIBRARY \
hipJitInputLibrary ///< Cuda Only Archive of Host Objects with embedded device code
#define HIPRTC_JIT_INPUT_NVVM \
hipJitInputNvvm ///< @deprecated CUDA only High Level intermediate code for LTO
#define HIPRTC_JIT_NUM_LEGACY_INPUT_TYPES \
hipJitNumLegacyInputTypes ///< Count of Legacy Input Types
#define HIPRTC_JIT_INPUT_LLVM_BITCODE \
hipJitInputLLVMBitcode ///< HIP Only LLVM Bitcode or IR assembly
#define HIPRTC_JIT_INPUT_LLVM_BUNDLED_BITCODE \
hipJitInputLLVMBundledBitcode ///< HIP Only LLVM Clang Bundled Code
#define HIPRTC_JIT_INPUT_LLVM_ARCHIVES_OF_BUNDLED_BITCODE \
hipJitInputLLVMArchivesOfBundledBitcode ///< HIP Only LLVM
///< Archives of
///< Bundled Bitcode
#define HIPRTC_JIT_INPUT_SPIRV hipJitInputSpirv ///< HIP Only SPIRV Code Object
#define HIPRTC_JIT_NUM_INPUT_TYPES hipJitNumInputTypes ///< Count of Input Types
/**
* @}
*/
/**
* hiprtc link state
*
*/
typedef struct ihiprtcLinkState* hiprtcLinkState;
/**
* @ingroup Runtime
*
* @brief Returns text string message to explain the error which occurred
*
* @param [in] result code to convert to string.
* @returns const char pointer to the NULL-terminated error string
*
* @warning In HIP, this function returns the name of the error,
* if the hiprtc result is defined, it will return "Invalid HIPRTC error code"
*
* @see hiprtcResult
*/
const char* hiprtcGetErrorString(hiprtcResult result);
/**
* @ingroup Runtime
* @brief Sets the parameters as major and minor version.
*
* @param [out] major HIP Runtime Compilation major version.
* @param [out] minor HIP Runtime Compilation minor version.
*
* @returns #HIPRTC_ERROR_INVALID_INPUT, #HIPRTC_SUCCESS
*
*/
hiprtcResult hiprtcVersion(int* major, int* minor);
/**
* hiprtc program
*
*/
typedef struct _hiprtcProgram* hiprtcProgram;
/**
* @ingroup Runtime
* @brief Adds the given name exprssion to the runtime compilation program.
*
* @param [in] prog runtime compilation program instance.
* @param [in] name_expression const char pointer to the name expression.
* @returns #HIPRTC_SUCCESS
*
* If const char pointer is NULL, it will return #HIPRTC_ERROR_INVALID_INPUT.
*
* @see hiprtcResult
*/
hiprtcResult hiprtcAddNameExpression(hiprtcProgram prog, const char* name_expression);
/**
* @ingroup Runtime
* @brief Compiles the given runtime compilation program.
*
* @param [in] prog runtime compilation program instance.
* @param [in] numOptions number of compiler options.
* @param [in] options compiler options as const array of strins.
* @returns #HIPRTC_SUCCESS
*
* If the compiler failed to build the runtime compilation program,
* it will return #HIPRTC_ERROR_COMPILATION.
*
* @see hiprtcResult
*/
hiprtcResult hiprtcCompileProgram(hiprtcProgram prog, int numOptions, const char* const* options);
/**
* @ingroup Runtime
* @brief Creates an instance of hiprtcProgram with the given input parameters,
* and sets the output hiprtcProgram prog with it.
*
* @param [in, out] prog runtime compilation program instance.
* @param [in] src const char pointer to the program source.
* @param [in] name const char pointer to the program name.
* @param [in] numHeaders number of headers.
* @param [in] headers array of strings pointing to headers.
* @param [in] includeNames array of strings pointing to names included in program source.
* @returns #HIPRTC_SUCCESS
*
* Any invalide input parameter, it will return #HIPRTC_ERROR_INVALID_INPUT
* or #HIPRTC_ERROR_INVALID_PROGRAM.
*
* If failed to create the program, it will return #HIPRTC_ERROR_PROGRAM_CREATION_FAILURE.
*
* @see hiprtcResult
*/
hiprtcResult hiprtcCreateProgram(hiprtcProgram* prog, const char* src, const char* name,
int numHeaders, const char* const* headers,
const char* const* includeNames);
/**
* @brief Destroys an instance of given hiprtcProgram.
* @ingroup Runtime
* @param [in] prog runtime compilation program instance.
* @returns #HIPRTC_SUCCESS
*
* If prog is NULL, it will return #HIPRTC_ERROR_INVALID_INPUT.
*
* @see hiprtcResult
*/
hiprtcResult hiprtcDestroyProgram(hiprtcProgram* prog);
/**
* @brief Gets the lowered (mangled) name from an instance of hiprtcProgram with the given input
* parameters, and sets the output lowered_name with it.
* @ingroup Runtime
* @param [in] prog runtime compilation program instance.
* @param [in] name_expression const char pointer to the name expression.
* @param [in, out] lowered_name const char array to the lowered (mangled) name.
* @returns #HIPRTC_SUCCESS
*
* If any invalide nullptr input parameters, it will return #HIPRTC_ERROR_INVALID_INPUT
*
* If name_expression is not found, it will return #HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID
*
* If failed to get lowered_name from the program, it will return #HIPRTC_ERROR_COMPILATION.
*
* @see hiprtcResult
*/
hiprtcResult hiprtcGetLoweredName(hiprtcProgram prog, const char* name_expression,
const char** lowered_name);
/**
* @brief Gets the log generated by the runtime compilation program instance.
* @ingroup Runtime
* @param [in] prog runtime compilation program instance.
* @param [out] log memory pointer to the generated log.
* @returns #HIPRTC_SUCCESS
*
* @see hiprtcResult
*/
hiprtcResult hiprtcGetProgramLog(hiprtcProgram prog, char* log);
/**
* @brief Gets the size of log generated by the runtime compilation program instance.
*
* @param [in] prog runtime compilation program instance.
* @param [out] logSizeRet size of generated log.
* @returns #HIPRTC_SUCCESS
*
* @see hiprtcResult
*/
hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram prog, size_t* logSizeRet);
/**
* @brief Gets the pointer of compilation binary by the runtime compilation program instance.
* @ingroup Runtime
* @param [in] prog runtime compilation program instance.
* @param [out] code char pointer to binary.
* @returns #HIPRTC_SUCCESS
*
* @see hiprtcResult
*/
hiprtcResult hiprtcGetCode(hiprtcProgram prog, char* code);
/**
* @brief Gets the size of compilation binary by the runtime compilation program instance.
* @ingroup Runtime
* @param [in] prog runtime compilation program instance.
* @param [out] codeSizeRet the size of binary.
* @returns #HIPRTC_SUCCESS
*
* @see hiprtcResult
*/
hiprtcResult hiprtcGetCodeSize(hiprtcProgram prog, size_t* codeSizeRet);
/**
* @brief Gets the pointer of compiled bitcode by the runtime compilation program instance.
*
* @param [in] prog runtime compilation program instance.
* @param [out] bitcode char pointer to bitcode.
* @return HIPRTC_SUCCESS
*
* @see hiprtcResult
*/
hiprtcResult hiprtcGetBitcode(hiprtcProgram prog, char* bitcode);
/**
* @brief Gets the size of compiled bitcode by the runtime compilation program instance.
* @ingroup Runtime
*
* @param [in] prog runtime compilation program instance.
* @param [out] bitcode_size the size of bitcode.
* @returns #HIPRTC_SUCCESS
*
* @see hiprtcResult
*/
hiprtcResult hiprtcGetBitcodeSize(hiprtcProgram prog, size_t* bitcode_size);
/**
* @brief Creates the link instance via hiprtc APIs.
* @ingroup Runtime
* @param [in] num_options Number of options
* @param [in] option_ptr Array of options
* @param [in] option_vals_pptr Array of option values cast to void*
* @param [out] hip_link_state_ptr hiprtc link state created upon success
*
* @returns #HIPRTC_SUCCESS, #HIPRTC_ERROR_INVALID_INPUT, #HIPRTC_ERROR_INVALID_OPTION
*
* @see hiprtcResult
*/
hiprtcResult hiprtcLinkCreate(unsigned int num_options, hiprtcJIT_option* option_ptr,
void** option_vals_pptr, hiprtcLinkState* hip_link_state_ptr);
/**
* @brief Adds a file with bit code to be linked with options
* @ingroup Runtime
* @param [in] hip_link_state hiprtc link state
* @param [in] input_type Type of the input data or bitcode
* @param [in] file_path Path to the input file where bitcode is present
* @param [in] num_options Size of the options
* @param [in] options_ptr Array of options applied to this input
* @param [in] option_values Array of option values cast to void*
*
* @returns #HIPRTC_SUCCESS
*
* If input values are invalid, it will
* @return #HIPRTC_ERROR_INVALID_INPUT
*
* @see hiprtcResult
*/
hiprtcResult hiprtcLinkAddFile(hiprtcLinkState hip_link_state, hiprtcJITInputType input_type,
const char* file_path, unsigned int num_options,
hiprtcJIT_option* options_ptr, void** option_values);
/**
* @brief Completes the linking of the given program.
* @ingroup Runtime
* @param [in] hip_link_state hiprtc link state
* @param [in] input_type Type of the input data or bitcode
* @param [in] image Input data which is null terminated
* @param [in] image_size Size of the input data
* @param [in] name Optional name for this input
* @param [in] num_options Size of the options
* @param [in] options_ptr Array of options applied to this input
* @param [in] option_values Array of option values cast to void*
*
* @returns #HIPRTC_SUCCESS, #HIPRTC_ERROR_INVALID_INPUT
*
* If adding the file fails, it will
* @return #HIPRTC_ERROR_PROGRAM_CREATION_FAILURE
*
* @see hiprtcResult
*/
hiprtcResult hiprtcLinkAddData(hiprtcLinkState hip_link_state, hiprtcJITInputType input_type,
void* image, size_t image_size, const char* name,
unsigned int num_options, hiprtcJIT_option* options_ptr,
void** option_values);
/**
* @brief Completes the linking of the given program.
* @ingroup Runtime
* @param [in] hip_link_state hiprtc link state
* @param [out] bin_out Upon success, points to the output binary
* @param [out] size_out Size of the binary is stored (optional)
*
* @returns #HIPRTC_SUCCESS
*
* If adding the data fails, it will
* @return #HIPRTC_ERROR_LINKING
*
* @see hiprtcResult
*/
hiprtcResult hiprtcLinkComplete(hiprtcLinkState hip_link_state, void** bin_out, size_t* size_out);
/**
* @brief Deletes the link instance via hiprtc APIs.
* @ingroup Runtime
* @param [in] hip_link_state link state instance
*
* @returns #HIPRTC_SUCCESS
*
* @see hiprtcResult
*/
hiprtcResult hiprtcLinkDestroy(hiprtcLinkState hip_link_state);
#if !defined(_WIN32)
#pragma GCC visibility pop
#endif
#ifdef __cplusplus
}
#endif /* __cplusplus */
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif

84
external/hip/hip/library_types.h vendored Normal file
View File

@@ -0,0 +1,84 @@
/*
Copyright (c) 2015 - 2023 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 HIP_INCLUDE_HIP_LIBRARY_TYPES_H
#define HIP_INCLUDE_HIP_LIBRARY_TYPES_H
#if !defined(__HIPCC_RTC__)
#include <hip/hip_common.h>
#endif
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
typedef enum hipDataType {
HIP_R_32F = 0,
HIP_R_64F = 1,
HIP_R_16F = 2,
HIP_R_8I = 3,
HIP_C_32F = 4,
HIP_C_64F = 5,
HIP_C_16F = 6,
HIP_C_8I = 7,
HIP_R_8U = 8,
HIP_C_8U = 9,
HIP_R_32I = 10,
HIP_C_32I = 11,
HIP_R_32U = 12,
HIP_C_32U = 13,
HIP_R_16BF = 14,
HIP_C_16BF = 15,
HIP_R_4I = 16,
HIP_C_4I = 17,
HIP_R_4U = 18,
HIP_C_4U = 19,
HIP_R_16I = 20,
HIP_C_16I = 21,
HIP_R_16U = 22,
HIP_C_16U = 23,
HIP_R_64I = 24,
HIP_C_64I = 25,
HIP_R_64U = 26,
HIP_C_64U = 27,
HIP_R_8F_E4M3 = 28,
HIP_R_8F_E5M2 = 29,
HIP_R_8F_UE8M0 = 30,
HIP_R_6F_E2M3 = 31,
HIP_R_6F_E3M2 = 32,
HIP_R_4F_E2M1 = 33,
// HIP specific Data Types
HIP_R_8F_E4M3_FNUZ = 1000,
HIP_R_8F_E5M2_FNUZ = 1001,
} hipDataType;
typedef enum hipLibraryPropertyType {
HIP_LIBRARY_MAJOR_VERSION,
HIP_LIBRARY_MINOR_VERSION,
HIP_LIBRARY_PATCH_LEVEL
} hipLibraryPropertyType;
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include "library_types.h"
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif

138
external/hip/hip/linker_types.h vendored Executable file
View File

@@ -0,0 +1,138 @@
/*
Copyright (c) 2015 - 2023 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 HIP_INCLUDE_HIP_LINKER_TYPES_H
#define HIP_INCLUDE_HIP_LINKER_TYPES_H
#if defined(__clang__)
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wreserved-identifier"
#pragma clang diagnostic ignored "-Wreserved-macro-identifier"
#endif
#if !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#elif defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
/**
* @defgroup LinkerTypes Jit Linker Data Types
* @{
* This section describes the Jit Linker data types.
*
*/
/**
* hipJitOption
*/
typedef enum hipJitOption {
hipJitOptionMaxRegisters = 0, ///< CUDA Only Maximum registers may be used in a thread,
///< passed to compiler
hipJitOptionThreadsPerBlock, ///< CUDA Only Number of thread per block
hipJitOptionWallTime, ///< CUDA Only Value for total wall clock time
hipJitOptionInfoLogBuffer, ///< CUDA Only Pointer to the buffer with logged information
hipJitOptionInfoLogBufferSizeBytes, ///< CUDA Only Size of the buffer in bytes for logged info
hipJitOptionErrorLogBuffer, ///< CUDA Only Pointer to the buffer with logged error(s)
hipJitOptionErrorLogBufferSizeBytes, ///< CUDA Only Size of the buffer in bytes for logged
///< error(s)
hipJitOptionOptimizationLevel, ///< Value of optimization level for generated codes, acceptable
///< options -O0, -O1, -O2, -O3
hipJitOptionTargetFromContext, ///< CUDA Only The target context, which is the default
hipJitOptionTarget, ///< CUDA Only JIT target
hipJitOptionFallbackStrategy, ///< CUDA Only Fallback strategy
hipJitOptionGenerateDebugInfo, ///< CUDA Only Generate debug information
hipJitOptionLogVerbose, ///< CUDA Only Generate log verbose
hipJitOptionGenerateLineInfo, ///< CUDA Only Generate line number information
hipJitOptionCacheMode, ///< CUDA Only Set cache mode
hipJitOptionSm3xOpt, ///< @deprecated CUDA Only New SM3X option.
hipJitOptionFastCompile, ///< CUDA Only Set fast compile
hipJitOptionGlobalSymbolNames, ///< CUDA Only Array of device symbol names to be relocated to the
///< host
hipJitOptionGlobalSymbolAddresses, ///< CUDA Only Array of host addresses to be relocated to the
///< device
hipJitOptionGlobalSymbolCount, ///< CUDA Only Number of symbol count.
hipJitOptionLto, ///< @deprecated CUDA Only Enable link-time optimization for device code
hipJitOptionFtz, ///< @deprecated CUDA Only Set single-precision denormals.
hipJitOptionPrecDiv, ///< @deprecated CUDA Only Set single-precision floating-point division
///< and reciprocals
hipJitOptionPrecSqrt, ///< @deprecated CUDA Only Set single-precision floating-point square root
hipJitOptionFma, ///< @deprecated CUDA Only Enable floating-point multiplies and
///< adds/subtracts operations
hipJitOptionPositionIndependentCode, ///< CUDA Only Generates Position Independent code
hipJitOptionMinCTAPerSM, ///< CUDA Only Hints to JIT compiler the minimum number of CTAs frin
///< kernel's grid to be mapped to SM
hipJitOptionMaxThreadsPerBlock, ///< CUDA only Maximum number of threads in a thread block
hipJitOptionOverrideDirectiveValues, ///< Cuda only Override Directive values
hipJitOptionNumOptions, ///< Number of options
hipJitOptionIRtoISAOptExt = 10000, ///< Hip Only Linker options to be passed on to compiler
hipJitOptionIRtoISAOptCountExt, ///< Hip Only Count of linker options to be passed on to compiler
} hipJitOption;
/**
* hipJitInputType
*/
typedef enum hipJitInputType {
hipJitInputCubin = 0, ///< Cuda only Input cubin
hipJitInputPtx, ///< Cuda only Input PTX
hipJitInputFatBinary, ///< Cuda Only Input FAT Binary
hipJitInputObject, ///< Cuda Only Host Object with embedded device code
hipJitInputLibrary, ///< Cuda Only Archive of Host Objects with embedded
///< device code
hipJitInputNvvm, ///< @deprecated Cuda only High Level intermediate
///< code for LTO
hipJitNumLegacyInputTypes, ///< Count of Legacy Input Types
hipJitInputLLVMBitcode = 100, ///< HIP Only LLVM Bitcode or IR assembly
hipJitInputLLVMBundledBitcode = 101, ///< HIP Only LLVM Clang Bundled Code
hipJitInputLLVMArchivesOfBundledBitcode = 102, ///< HIP Only LLVM Archive of Bundled Bitcode
hipJitInputSpirv = 103, ///< HIP Only SPIRV Code Object
hipJitNumInputTypes = 10 ///< Count of Input Types
} hipJitInputType;
/**
* hipJitCacheMode
*/
typedef enum hipJitCacheMode {
hipJitCacheOptionNone = 0,
hipJitCacheOptionCG,
hipJitCacheOptionCA
} hipJitCacheMode;
/**
* hipJitFallback
*/
typedef enum hipJitFallback {
hipJitPreferPTX = 0,
hipJitPreferBinary,
} hipJitFallback;
typedef enum hipLibraryOption_e {
hipLibraryHostUniversalFunctionAndDataTable = 0,
hipLibraryBinaryIsPreserved = 1
} hipLibraryOption;
// doxygen end LinkerTypes
/**
* @}
*/
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif // HIP_INCLUDE_HIP_LINKER_TYPES_H

42
external/hip/hip/math_functions.h vendored Normal file
View File

@@ -0,0 +1,42 @@
/*
Copyright (c) 2015 - 2023 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 HIP_INCLUDE_HIP_MATH_FUNCTIONS_H
#define HIP_INCLUDE_HIP_MATH_FUNCTIONS_H
// Some standard header files, these are included by hc.hpp and so want to make them avail on both
// paths to provide a consistent include env and avoid "missing symbol" errors that only appears
// on NVCC path:
#if !defined(__HIPCC_RTC__)
#include <hip/hip_common.h>
#endif
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/amd_detail/amd_math_functions.h>
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
// #include <hip/nvidia_detail/math_functions.h>
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif

65
external/hip/hip/surface_types.h vendored Normal file
View File

@@ -0,0 +1,65 @@
/*
Copyright (c) 2022 - 2023 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.
*/
/**
* @file surface_types.h
* @brief Defines surface types for HIP runtime.
*/
#ifndef HIP_INCLUDE_HIP_SURFACE_TYPES_H
#define HIP_INCLUDE_HIP_SURFACE_TYPES_H
#if defined(__clang__)
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wreserved-identifier"
#endif
#if !defined(__HIPCC_RTC__)
#include <hip/driver_types.h>
#endif
/**
* An opaque value that represents a hip surface object
*/
struct __hip_surface;
typedef struct __hip_surface* hipSurfaceObject_t;
/**
* hip surface reference
*/
struct surfaceReference {
hipSurfaceObject_t surfaceObject;
};
/**
* hip surface boundary modes
*/
enum hipSurfaceBoundaryMode {
hipBoundaryModeZero = 0,
hipBoundaryModeTrap = 1,
hipBoundaryModeClamp = 2
};
#if defined(__clang__)
#pragma clang diagnostic pop
#endif
#endif /* !HIP_INCLUDE_HIP_SURFACE_TYPES_H */

193
external/hip/hip/texture_types.h vendored Normal file
View File

@@ -0,0 +1,193 @@
/*
Copyright (c) 2015 - 2023 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 HIP_INCLUDE_HIP_TEXTURE_TYPES_H
#define HIP_INCLUDE_HIP_TEXTURE_TYPES_H
#if defined(__clang__)
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wreserved-identifier"
#pragma clang diagnostic ignored "-Wreserved-macro-identifier"
#pragma clang diagnostic ignored "-Wc++98-compat"
#endif
#if !defined(__HIPCC_RTC__)
#include <hip/hip_common.h>
#endif
#if !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include "texture_types.h"
#elif defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
/*******************************************************************************
* *
* *
* *
*******************************************************************************/
#if !defined(__HIPCC_RTC__)
#include <hip/channel_descriptor.h>
#include <hip/driver_types.h>
#endif // !defined(__HIPCC_RTC__)
#define hipTextureType1D 0x01
#define hipTextureType2D 0x02
#define hipTextureType3D 0x03
#define hipTextureTypeCubemap 0x0C
#define hipTextureType1DLayered 0xF1
#define hipTextureType2DLayered 0xF2
#define hipTextureTypeCubemapLayered 0xFC
/**
* Should be same as HSA_IMAGE_OBJECT_SIZE_DWORD/HSA_SAMPLER_OBJECT_SIZE_DWORD
*/
#define HIP_IMAGE_OBJECT_SIZE_DWORD 12
#define HIP_SAMPLER_OBJECT_SIZE_DWORD 8
#define HIP_SAMPLER_OBJECT_OFFSET_DWORD HIP_IMAGE_OBJECT_SIZE_DWORD
#define HIP_TEXTURE_OBJECT_SIZE_DWORD (HIP_IMAGE_OBJECT_SIZE_DWORD + HIP_SAMPLER_OBJECT_SIZE_DWORD)
/**
* An opaque value that represents a hip texture object
*/
struct __hip_texture;
typedef struct __hip_texture* hipTextureObject_t;
/**
* hip texture address modes
*/
enum hipTextureAddressMode {
hipAddressModeWrap = 0,
hipAddressModeClamp = 1,
hipAddressModeMirror = 2,
hipAddressModeBorder = 3
};
/**
* hip texture filter modes
*/
enum hipTextureFilterMode { hipFilterModePoint = 0, hipFilterModeLinear = 1 };
/**
* hip texture read modes
*/
enum hipTextureReadMode { hipReadModeElementType = 0, hipReadModeNormalizedFloat = 1 };
/**
* hip texture reference
*/
typedef struct textureReference {
int normalized;
enum hipTextureReadMode readMode; // used only for driver API's
enum hipTextureFilterMode filterMode;
enum hipTextureAddressMode addressMode[3]; // Texture address mode for up to 3 dimensions
struct hipChannelFormatDesc channelDesc;
int sRGB; // Perform sRGB->linear conversion during texture read
unsigned int maxAnisotropy; // Limit to the anisotropy ratio
enum hipTextureFilterMode mipmapFilterMode;
float mipmapLevelBias;
float minMipmapLevelClamp;
float maxMipmapLevelClamp;
hipTextureObject_t textureObject;
int numChannels;
enum hipArray_Format format;
} textureReference;
/**
* hip texture descriptor
*/
typedef struct hipTextureDesc {
enum hipTextureAddressMode addressMode[3]; // Texture address mode for up to 3 dimensions
enum hipTextureFilterMode filterMode;
enum hipTextureReadMode readMode;
int sRGB; // Perform sRGB->linear conversion during texture read
float borderColor[4];
int normalizedCoords;
unsigned int maxAnisotropy;
enum hipTextureFilterMode mipmapFilterMode;
float mipmapLevelBias;
float minMipmapLevelClamp;
float maxMipmapLevelClamp;
} hipTextureDesc;
#if __cplusplus
/*******************************************************************************
* *
* *
* *
*******************************************************************************/
#if __HIP__
#define __HIP_TEXTURE_ATTRIB __attribute__((device_builtin_texture_type))
#else
#define __HIP_TEXTURE_ATTRIB
#endif
typedef textureReference* hipTexRef;
template <class T, int texType = hipTextureType1D,
enum hipTextureReadMode mode = hipReadModeElementType>
struct __HIP_TEXTURE_ATTRIB texture : public textureReference {
texture(int norm = 0, enum hipTextureFilterMode fMode = hipFilterModePoint,
enum hipTextureAddressMode aMode = hipAddressModeClamp) {
normalized = norm;
readMode = mode;
filterMode = fMode;
addressMode[0] = aMode;
addressMode[1] = aMode;
addressMode[2] = aMode;
channelDesc = hipCreateChannelDesc<T>();
sRGB = 0;
textureObject = nullptr;
maxAnisotropy = 0;
mipmapLevelBias = 0;
minMipmapLevelClamp = 0;
maxMipmapLevelClamp = 0;
}
texture(int norm, enum hipTextureFilterMode fMode, enum hipTextureAddressMode aMode,
struct hipChannelFormatDesc desc) {
normalized = norm;
readMode = mode;
filterMode = fMode;
addressMode[0] = aMode;
addressMode[1] = aMode;
addressMode[2] = aMode;
channelDesc = desc;
sRGB = 0;
textureObject = nullptr;
maxAnisotropy = 0;
mipmapLevelBias = 0;
minMipmapLevelClamp = 0;
maxMipmapLevelClamp = 0;
}
};
#endif /* __cplusplus */
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#if defined(__clang__)
#pragma clang diagnostic pop
#endif
#endif

36
external/hip/hip_bf16.h vendored Normal file
View File

@@ -0,0 +1,36 @@
/*
Copyright (c) 2023 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 HIP_INCLUDE_HIP_HIP_BF16_H
#define HIP_INCLUDE_HIP_HIP_BF16_H
#include <hip/hip_common.h>
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/amd_detail/amd_hip_bf16.h>
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/nvidia_detail/nvidia_hip_bf16.h>
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif // HIP_INCLUDE_HIP_HIP_BF16_H

44
external/hip/hip_bfloat16.h vendored Normal file
View File

@@ -0,0 +1,44 @@
/**
* MIT License
*
* Copyright (c) 2019 - 2022 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.
*/
/*!\file
* \brief hip_bfloat16.h provides struct for hip_bfloat16 typedef
*/
#ifndef _HIP_BFLOAT16_H_
#define _HIP_BFLOAT16_H_
#if !defined(__HIPCC_RTC__)
#include <hip/hip_common.h>
#endif
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/amd_detail/amd_hip_bfloat16.h>
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#warning "hip_bfloat16.h is not supported on nvidia platform"
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif // _HIP_BFLOAT16_H_

100
external/hip/hip_common.h vendored Normal file
View File

@@ -0,0 +1,100 @@
/*
Copyright (c) 2015 - 2023 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 HIP_INCLUDE_HIP_HIP_COMMON_H
#define HIP_INCLUDE_HIP_HIP_COMMON_H
#if defined(__clang__)
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wreserved-macro-identifier"
#endif
// Common code included at start of every hip file.
// Auto enable __HIP_PLATFORM_AMD__ if compiling on AMD platform
// Other compiler (GCC,ICC,etc) need to set one of these macros explicitly
#if defined(__clang__) && defined(__HIP__)
#ifndef __HIP_PLATFORM_AMD__
#define __HIP_PLATFORM_AMD__
#endif
#endif // defined(__clang__) && defined(__HIP__)
// Auto enable __HIP_PLATFORM_NVIDIA__ if compiling with NVIDIA platform
#if defined(__NVCC__) || (defined(__clang__) && defined(__CUDA__) && !defined(__HIP__))
#ifndef __HIP_PLATFORM_NVIDIA__
#define __HIP_PLATFORM_NVIDIA__
#endif
#ifdef __CUDACC__
#define __HIPCC__
#endif
#endif //__NVCC__
// Auto enable __HIP_DEVICE_COMPILE__ if compiled in HCC or NVCC device path
#if (defined(__HCC_ACCELERATOR__) && __HCC_ACCELERATOR__ != 0) || \
(defined(__CUDA_ARCH__) && __CUDA_ARCH__ != 0)
#define __HIP_DEVICE_COMPILE__ 1
#endif
#ifdef __GNUC__
#define HIP_PUBLIC_API __attribute__((visibility("default")))
#define HIP_INTERNAL_EXPORTED_API __attribute__((visibility("default")))
#else
#define HIP_PUBLIC_API
#define HIP_INTERNAL_EXPORTED_API
#endif
#if __HIP_DEVICE_COMPILE__ == 0
// 32-bit Atomics
#define __HIP_ARCH_HAS_GLOBAL_INT32_ATOMICS__ (0)
#define __HIP_ARCH_HAS_GLOBAL_FLOAT_ATOMIC_EXCH__ (0)
#define __HIP_ARCH_HAS_SHARED_INT32_ATOMICS__ (0)
#define __HIP_ARCH_HAS_SHARED_FLOAT_ATOMIC_EXCH__ (0)
#define __HIP_ARCH_HAS_FLOAT_ATOMIC_ADD__ (0)
// 64-bit Atomics
#define __HIP_ARCH_HAS_GLOBAL_INT64_ATOMICS__ (0)
#define __HIP_ARCH_HAS_SHARED_INT64_ATOMICS__ (0)
// Doubles
#define __HIP_ARCH_HAS_DOUBLES__ (0)
// Warp cross-lane operations
#define __HIP_ARCH_HAS_WARP_VOTE__ (0)
#define __HIP_ARCH_HAS_WARP_BALLOT__ (0)
#define __HIP_ARCH_HAS_WARP_SHUFFLE__ (0)
#define __HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__ (0)
// Sync
#define __HIP_ARCH_HAS_THREAD_FENCE_SYSTEM__ (0)
#define __HIP_ARCH_HAS_SYNC_THREAD_EXT__ (0)
// Misc
#define __HIP_ARCH_HAS_SURFACE_FUNCS__ (0)
#define __HIP_ARCH_HAS_3DGRID__ (0)
#define __HIP_ARCH_HAS_DYNAMIC_PARALLEL__ (0)
#endif
#if defined(__clang__)
#pragma clang diagnostic pop
#endif
#endif

38
external/hip/hip_complex.h vendored Normal file
View File

@@ -0,0 +1,38 @@
/*
Copyright (c) 2015 - 2023 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 HIP_INCLUDE_HIP_HIP_COMPLEX_H
#define HIP_INCLUDE_HIP_HIP_COMPLEX_H
#if !defined(__HIPCC_RTC__)
#include <hip/hip_common.h>
#endif
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/amd_detail/amd_hip_complex.h>
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/nvidia_detail/nvidia_hip_complex.h>
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif

46
external/hip/hip_cooperative_groups.h vendored Normal file
View File

@@ -0,0 +1,46 @@
/*
Copyright (c) 2015 - 2021 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.
*/
/**
* @file hip_cooperative_groups.h
*
* @brief Defines new types and device API wrappers for `Cooperative Group`
* feature.
*/
#ifndef HIP_INCLUDE_HIP_HIP_COOPERATIVE_GROUP_H
#define HIP_INCLUDE_HIP_HIP_COOPERATIVE_GROUP_H
#include <hip/hip_version.h>
#include <hip/hip_common.h>
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#if __cplusplus && defined(__clang__) && defined(__HIP__)
#include <hip/amd_detail/amd_hip_cooperative_groups.h>
#endif
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/nvidia_detail/nvidia_hip_cooperative_groups.h>
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif // HIP_INCLUDE_HIP_HIP_COOPERATIVE_GROUP_H

119
external/hip/hip_deprecated.h vendored Normal file
View File

@@ -0,0 +1,119 @@
/*
* Copyright (C) Advanced Micro Devices, Inc.
*
* 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 COPYRIGHT HOLDER(S) 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.
*/
#pragma once
// This file will add older hip functions used in the versioning system
// Find the deprecated functions and structs in hip_device.cpp
// This struct is also kept in hip_device.cpp
typedef struct hipDeviceProp_tR0000 {
char name[256]; ///< Device name.
size_t totalGlobalMem; ///< Size of global memory region (in bytes).
size_t sharedMemPerBlock; ///< Size of shared memory region (in bytes).
int regsPerBlock; ///< Registers per block.
int warpSize; ///< Warp size.
int maxThreadsPerBlock; ///< Max work items per work group or workgroup max size.
int maxThreadsDim[3]; ///< Max number of threads in each dimension (XYZ) of a block.
int maxGridSize[3]; ///< Max grid dimensions (XYZ).
int clockRate; ///< Max clock frequency of the multiProcessors in khz.
int memoryClockRate; ///< Max global memory clock frequency in khz.
int memoryBusWidth; ///< Global memory bus width in bits.
size_t totalConstMem; ///< Size of shared memory region (in bytes).
int major; ///< Major compute capability. On HCC, this is an approximation and features may
///< differ from CUDA CC. See the arch feature flags for portable ways to query
///< feature caps.
int minor; ///< Minor compute capability. On HCC, this is an approximation and features may
///< differ from CUDA CC. See the arch feature flags for portable ways to query
///< feature caps.
int multiProcessorCount; ///< Number of multi-processors. When the GPU works in Compute
///< Unit (CU) mode, this value equals the number of CUs;
///< when in Workgroup Processor (WGP) mode, this value equels
///< half of CUs, because a single WGP contains two CUs.
int l2CacheSize; ///< L2 cache size.
int maxThreadsPerMultiProcessor; ///< Maximum resident threads per multi-processor.
int computeMode; ///< Compute mode.
int clockInstructionRate; ///< Frequency in khz of the timer used by the device-side "clock*"
///< instructions. New for HIP.
hipDeviceArch_t arch; ///< Architectural feature flags. New for HIP.
int concurrentKernels; ///< Device can possibly execute multiple kernels concurrently.
int pciDomainID; ///< PCI Domain ID
int pciBusID; ///< PCI Bus ID.
int pciDeviceID; ///< PCI Device ID.
size_t maxSharedMemoryPerMultiProcessor; ///< Maximum Shared Memory Per Multiprocessor.
int isMultiGpuBoard; ///< 1 if device is on a multi-GPU board, 0 if not.
int canMapHostMemory; ///< Check whether HIP can map host memory
int gcnArch; ///< DEPRECATED: use gcnArchName instead
char gcnArchName[256]; ///< AMD GCN Arch Name.
int integrated; ///< APU vs dGPU
int cooperativeLaunch; ///< HIP device supports cooperative launch
int cooperativeMultiDeviceLaunch; ///< HIP device supports cooperative launch on multiple
///< devices
int maxTexture1DLinear; ///< Maximum size for 1D textures bound to linear memory
int maxTexture1D; ///< Maximum number of elements in 1D images
int maxTexture2D[2]; ///< Maximum dimensions (width, height) of 2D images, in image elements
int maxTexture3D[3]; ///< Maximum dimensions (width, height, depth) of 3D images, in image
///< elements
unsigned int* hdpMemFlushCntl; ///< Addres of HDP_MEM_COHERENCY_FLUSH_CNTL register
unsigned int* hdpRegFlushCntl; ///< Addres of HDP_REG_COHERENCY_FLUSH_CNTL register
size_t memPitch; ///< Maximum pitch in bytes allowed by memory copies
size_t textureAlignment; ///< Alignment requirement for textures
size_t texturePitchAlignment; ///< Pitch alignment requirement for texture references bound to
///< pitched memory
int kernelExecTimeoutEnabled; ///< Run time limit for kernels executed on the device
int ECCEnabled; ///< Device has ECC support enabled
int tccDriver; ///< 1:If device is Tesla device using TCC driver, else 0
int cooperativeMultiDeviceUnmatchedFunc; ///< HIP device supports cooperative launch on
///< multiple
/// devices with unmatched functions
int cooperativeMultiDeviceUnmatchedGridDim; ///< HIP device supports cooperative launch on
///< multiple
/// devices with unmatched grid dimensions
int cooperativeMultiDeviceUnmatchedBlockDim; ///< HIP device supports cooperative launch on
///< multiple
/// devices with unmatched block dimensions
int cooperativeMultiDeviceUnmatchedSharedMem; ///< HIP device supports cooperative launch on
///< multiple
/// devices with unmatched shared memories
int isLargeBar; ///< 1: if it is a large PCI bar device, else 0
int asicRevision; ///< Revision of the GPU in this device
int managedMemory; ///< Device supports allocating managed memory on this system
int directManagedMemAccessFromHost; ///< Host can directly access managed memory on the device
///< without migration
int concurrentManagedAccess; ///< Device can coherently access managed memory concurrently with
///< the CPU
int pageableMemoryAccess; ///< Device supports coherently accessing pageable memory
///< without calling hipHostRegister on it
int pageableMemoryAccessUsesHostPageTables; ///< Device accesses pageable memory via the host's
///< page tables
} hipDeviceProp_tR0000;
#ifdef __cplusplus
extern "C" {
#endif
hipError_t hipGetDevicePropertiesR0000(hipDeviceProp_tR0000* prop, int device);
hipError_t hipChooseDeviceR0000(int* device, const hipDeviceProp_tR0000* prop);
#ifdef __cplusplus
}
#endif

162
external/hip/hip_ext.h vendored Normal file
View File

@@ -0,0 +1,162 @@
/*
Copyright (c) 2015 - 2021 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 HIP_INCLUDE_HIP_HIP_EXT_H
#define HIP_INCLUDE_HIP_HIP_EXT_H
#include "hip/hip_runtime.h"
#if defined(__cplusplus)
#include <tuple>
#include <type_traits>
#endif
/** @addtogroup Execution Execution Control
* @{
*/
/**
* @brief Launches kernel with parameters and shared memory on stream with arguments passed
* to kernel params or extra arguments.
*
* @param [in] f Kernel to launch.
* @param [in] globalWorkSizeX X grid dimension specified in work-items.
* @param [in] globalWorkSizeY Y grid dimension specified in work-items.
* @param [in] globalWorkSizeZ Z grid dimension specified in work-items.
* @param [in] localWorkSizeX X block dimension specified in work-items.
* @param [in] localWorkSizeY Y block dimension specified in work-items.
* @param [in] localWorkSizeZ Z block dimension specified in work-items.
* @param [in] sharedMemBytes Amount of dynamic shared memory to allocate for this kernel.
* HIP-Clang compiler provides support for extern shared declarations.
* @param [in] hStream Stream where the kernel should be dispatched.
* May be 0, in which case the default stream is used with associated synchronization rules.
* @param [in] kernelParams pointer to kernel parameters.
* @param [in] extra Pointer to kernel arguments. These are passed directly to the kernel and
* must be in the memory layout and alignment expected by the kernel.
* All passed arguments must be naturally aligned according to their type. The memory address of
* each argument should be a multiple of its size in bytes. Please refer to
* hip_porting_driver_api.md for sample usage.
* @param [in] startEvent If non-null, specified event will be updated to track the start time of
* the kernel launch. The event must be created before calling this API.
* @param [in] stopEvent If non-null, specified event will be updated to track the stop time of
* the kernel launch. The event must be created before calling this API.
* @param [in] flags The value of hipExtAnyOrderLaunch, signifies if kernel can be
* launched in any order.
* @returns #hipSuccess, #hipInvalidDeviceId, #hipErrorNotInitialized, #hipErrorInvalidValue.
*
* HIP/ROCm actually updates the start event when the associated kernel completes.
* Currently, timing between startEvent and stopEvent does not include the time it takes to perform
* a system scope release/cache flush - only the time it takes to issues writes to cache.
*
* @note For this HIP API, the flag 'hipExtAnyOrderLaunch' is not supported on AMD GFX9xx boards.
*
*/
HIP_PUBLIC_API
extern "C" hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
uint32_t localWorkSizeX, uint32_t localWorkSizeY,
uint32_t localWorkSizeZ, size_t sharedMemBytes,
hipStream_t hStream, void** kernelParams,
void** extra, hipEvent_t startEvent __dparm(NULL),
hipEvent_t stopEvent __dparm(NULL),
uint32_t flags __dparm(0));
/**
* @brief This HIP API is deprecated, please use hipExtModuleLaunchKernel() instead.
*
*/
HIP_DEPRECATED("use hipExtModuleLaunchKernel instead")
HIP_PUBLIC_API
extern "C" hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
uint32_t localWorkSizeX, uint32_t localWorkSizeY,
uint32_t localWorkSizeZ, size_t sharedMemBytes,
hipStream_t hStream, void** kernelParams,
void** extra, hipEvent_t startEvent __dparm(NULL),
hipEvent_t stopEvent __dparm(NULL));
#if defined(__cplusplus)
/**
* @brief Launches kernel from the pointer address, with arguments and shared memory on stream.
*
* @param [in] function_address pointer to the Kernel to launch.
* @param [in] numBlocks number of blocks.
* @param [in] dimBlocks dimension of a block.
* @param [in] args pointer to kernel arguments.
* @param [in] sharedMemBytes Amount of dynamic shared memory to allocate for this kernel.
* HIP-Clang compiler provides support for extern shared declarations.
* @param [in] stream Stream where the kernel should be dispatched.
* May be 0, in which case the default stream is used with associated synchronization rules.
* @param [in] startEvent If non-null, specified event will be updated to track the start time of
* the kernel launch. The event must be created before calling this API.
* @param [in] stopEvent If non-null, specified event will be updated to track the stop time of
* the kernel launch. The event must be created before calling this API.
* @param [in] flags The value of hipExtAnyOrderLaunch, signifies if kernel can be
* launched in any order.
* @returns #hipSuccess, #hipInvalidDeviceId, #hipErrorNotInitialized, #hipErrorInvalidValue.
*
*/
extern "C" hipError_t hipExtLaunchKernel(const void* function_address, dim3 numBlocks,
dim3 dimBlocks, void** args, size_t sharedMemBytes,
hipStream_t stream, hipEvent_t startEvent,
hipEvent_t stopEvent, int flags);
/**
* @brief Launches kernel with dimention parameters and shared memory on stream with templated
* kernel and arguments.
*
* @param [in] kernel Kernel to launch.
* @param [in] numBlocks const number of blocks.
* @param [in] dimBlocks const dimension of a block.
* @param [in] sharedMemBytes Amount of dynamic shared memory to allocate for this kernel.
* HIP-Clang compiler provides support for extern shared declarations.
* @param [in] stream Stream where the kernel should be dispatched.
* May be 0, in which case the default stream is used with associated synchronization rules.
* @param [in] startEvent If non-null, specified event will be updated to track the start time of
* the kernel launch. The event must be created before calling this API.
* @param [in] stopEvent If non-null, specified event will be updated to track the stop time of
* the kernel launch. The event must be created before calling this API.
* @param [in] flags The value of hipExtAnyOrderLaunch, signifies if kernel can be
* launched in any order.
* @param [in] args templated kernel arguments.
*
*/
template <typename... Args, typename F = void (*)(Args...)>
inline void hipExtLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
std::uint32_t sharedMemBytes, hipStream_t stream,
hipEvent_t startEvent, hipEvent_t stopEvent, std::uint32_t flags,
Args... args) {
constexpr size_t count = sizeof...(Args);
auto tup_ = std::tuple<Args...>{args...};
auto tup = validateArgsCountType(kernel, tup_);
void* _Args[count];
pArgs<0>(tup, _Args);
auto k = reinterpret_cast<void*>(kernel);
hipExtLaunchKernel(k, numBlocks, dimBlocks, _Args, sharedMemBytes, stream, startEvent, stopEvent,
(int)flags);
}
#endif // defined(__cplusplus)
// doxygen end AMD-specific features
/**
* @}
*/
#endif // #iidef HIP_INCLUDE_HIP_HIP_EXT_H

31
external/hip/hip_ext_ocp.h vendored Normal file
View File

@@ -0,0 +1,31 @@
/*
Copyright (c) 2024 - present 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 HIP_INCLUDE_HIP_HIP_EXT_OCP_H
#define HIP_INCLUDE_HIP_HIP_EXT_OCP_H
#include <hip/amd_detail/amd_hip_ocp_types.h>
#include <hip/amd_detail/amd_hip_ocp_fp.hpp>
#include <hip/amd_detail/amd_hip_ocp_fp_cxx.hpp>
#endif // HIP_INCLUDE_HIP_HIP_EXT_OCP_H

44
external/hip/hip_fp16.h vendored Normal file
View File

@@ -0,0 +1,44 @@
/*
Copyright (c) 2015 - 2021 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 HIP_INCLUDE_HIP_HIP_FP16_H
#define HIP_INCLUDE_HIP_HIP_FP16_H
#include <hip/hip_common.h>
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/amd_detail/amd_hip_fp16.h>
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#define HIPRT_INF_FP16 CUDART_INF_FP16
#define HIPRT_MAX_NORMAL_FP16 CUDART_MAX_NORMAL_FP16
#define HIPRT_MIN_DENORM_FP16 CUDART_MIN_DENORM_FP16
#define HIPRT_NAN_FP16 CUDART_NAN_FP16
#define HIPRT_NEG_ZERO_FP16 CUDART_NEG_ZERO_FP16
#define HIPRT_ONE_FP16 CUDART_ONE_FP16
#define HIPRT_ZERO_FP16 CUDART_ZERO_FP16
#include "cuda_fp16.h"
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif

32
external/hip/hip_fp4.h vendored Normal file
View File

@@ -0,0 +1,32 @@
/*
Copyright (c) 2025 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 HIP_INCLUDE_HIP_HIP_FP4_H
#define HIP_INCLUDE_HIP_HIP_FP4_H
#include <hip/hip_common.h>
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/amd_detail/amd_hip_fp4.h>
#endif
#endif // HIP_INCLUDE_HIP_HIP_FP4_H

32
external/hip/hip_fp6.h vendored Normal file
View File

@@ -0,0 +1,32 @@
/*
Copyright (c) 2025 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 HIP_INCLUDE_HIP_HIP_FP6_H
#define HIP_INCLUDE_HIP_HIP_FP6_H
#include <hip/hip_common.h>
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/amd_detail/amd_hip_fp6.h>
#endif
#endif // HIP_INCLUDE_HIP_HIP_FP6_H

33
external/hip/hip_fp8.h vendored Normal file
View File

@@ -0,0 +1,33 @@
/*
Copyright (c) 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 HIP_INCLUDE_HIP_HIP_FP8_H
#define HIP_INCLUDE_HIP_HIP_FP8_H
#include <hip/hip_common.h>
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
// We only have fnuz defs for now, which are not supported by other platforms
#include <hip/amd_detail/amd_hip_fp8.h>
#endif
#endif // HIP_INCLUDE_HIP_HIP_FP8_H

32
external/hip/hip_gl_interop.h vendored Normal file
View File

@@ -0,0 +1,32 @@
/*
Copyright (c) 2023 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 HIP_GL_INTEROP_H
#define HIP_GL_INTEROP_H
#include <hip/hip_common.h>
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#include "hip/amd_detail/amd_hip_gl_interop.h"
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include "hip/nvidia_detail/nvidia_hip_gl_interop.h"
#endif
#endif

24
external/hip/hip_hcc.h vendored Normal file
View File

@@ -0,0 +1,24 @@
/*
Copyright (c) 2015 - 2021 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 HIP_INCLUDE_HIP_HIP_HCC_H
#define HIP_INCLUDE_HIP_HIP_HCC_H
#warning "hip/hip_hcc.h is deprecated, please use hip/hip_ext.h"
#include "hip/hip_ext.h"
#endif // #ifdef HIP_INCLUDE_HIP_HIP_HCC_H

36
external/hip/hip_math_constants.h vendored Normal file
View File

@@ -0,0 +1,36 @@
/*
Copyright (c) 2015 - 2022 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 HIP_MATH_CONSTANTS_H
#define HIP_MATH_CONSTANTS_H
#if !defined(__HIPCC_RTC__)
#include <hip/hip_common.h>
#endif
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#include "hip/amd_detail/amd_hip_math_constants.h"
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include "hip/nvidia_detail/nvidia_hip_math_constants.h"
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif

27
external/hip/hip_profile.h vendored Normal file
View File

@@ -0,0 +1,27 @@
/*
Copyright (c) 2015 - 2021 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 HIP_INCLUDE_HIP_HIP_PROFILE_H
#define HIP_INCLUDE_HIP_HIP_PROFILE_H
#define HIP_SCOPED_MARKER(markerName, group)
#define HIP_BEGIN_MARKER(markerName, group)
#define HIP_END_MARKER()
#endif

70
external/hip/hip_runtime.h vendored Normal file
View File

@@ -0,0 +1,70 @@
/*
Copyright (c) 2015 - 2025 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.
*/
//! HIP = Heterogeneous-compute Interface for Portability
//!
//! Define a extremely thin runtime layer that allows source code to be compiled unmodified
//! through either AMD CLANG or NVCC. Key features tend to be in the spirit
//! and terminology of CUDA, but with a portable path to other accelerators as well:
//
//! Both paths support rich C++ features including classes, templates, lambdas, etc.
//! Runtime API is C
//! Memory management is based on pure pointers and resembles malloc/free/copy.
//
//! hip_runtime.h : includes everything in hip_api.h, plus math builtins and kernel launch
//! macros. hip_runtime_api.h : Defines HIP API. This is a C header file and does not use any C++
//! features.
#ifndef HIP_INCLUDE_HIP_HIP_RUNTIME_H
#define HIP_INCLUDE_HIP_HIP_RUNTIME_H
#if !defined(__HIPCC_RTC__)
// Some standard header files, these are included by hc.hpp and so want to make them avail on both
// paths to provide a consistent include env and avoid "missing symbol" errors that only appears
// on NVCC path:
#if __cplusplus
#include <cstdint>
#include <cstdlib>
#else
#include <stdint.h>
#include <stdlib.h>
#endif // __cplusplus
#endif // !defined(__HIPCC_RTC__)
#include <hip/hip_version.h>
#include <hip/hip_common.h>
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/amd_detail/amd_hip_runtime.h>
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/nvidia_detail/nvidia_hip_runtime.h>
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#if !defined(__HIPCC_RTC__)
#include <hip/hip_runtime_api.h>
#include <hip/library_types.h>
#endif // !defined(__HIPCC_RTC__)
#include <hip/hip_vector_types.h>
#endif

10431
external/hip/hip_runtime_api.h vendored Normal file
View File

File diff suppressed because it is too large Load Diff

29
external/hip/hip_texture_types.h vendored Normal file
View File

@@ -0,0 +1,29 @@
/*
Copyright (c) 2015 - 2021 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 HIP_INCLUDE_HIP_HIP_TEXTURE_TYPES_H
#define HIP_INCLUDE_HIP_HIP_TEXTURE_TYPES_H
#include <hip/texture_types.h>
#endif

41
external/hip/hip_vector_types.h vendored Normal file
View File

@@ -0,0 +1,41 @@
/*
Copyright (c) 2015 - 2021 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.
*/
//! hip_vector_types.h : Defines the HIP vector types.
#ifndef HIP_INCLUDE_HIP_HIP_VECTOR_TYPES_H
#define HIP_INCLUDE_HIP_HIP_VECTOR_TYPES_H
#include <hip/hip_common.h>
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#if __cplusplus
#include <hip/amd_detail/amd_hip_vector_types.h>
#endif
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include <vector_types.h>
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif

473
external/hip/hiprtc.h vendored Normal file
View File

@@ -0,0 +1,473 @@
/*
Copyright (c) 2015 - 2023 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.
*/
#pragma once
#include <hip/hip_common.h>
#include <hip/linker_types.h>
#if !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/nvidia_detail/nvidia_hiprtc.h>
#elif defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#ifdef __cplusplus
#include <cstdlib>
#else
#include <stdlib.h>
#endif
#ifdef __cplusplus
extern "C" {
#endif /* __cplusplus */
#if !defined(_WIN32)
#pragma GCC visibility push(default)
#endif
/**
*
* @addtogroup GlobalDefs
* @{
*
*/
/**
* hiprtc error code
*/
typedef enum hiprtcResult {
HIPRTC_SUCCESS = 0, ///< Success
HIPRTC_ERROR_OUT_OF_MEMORY = 1, ///< Out of memory
HIPRTC_ERROR_PROGRAM_CREATION_FAILURE = 2, ///< Failed to create program
HIPRTC_ERROR_INVALID_INPUT = 3, ///< Invalid input
HIPRTC_ERROR_INVALID_PROGRAM = 4, ///< Invalid program
HIPRTC_ERROR_INVALID_OPTION = 5, ///< Invalid option
HIPRTC_ERROR_COMPILATION = 6, ///< Compilation error
HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE = 7, ///< Failed in builtin operation
HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = 8, ///< No name expression after compilation
HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = 9, ///< No lowered names before compilation
HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID = 10, ///< Invalid name expression
HIPRTC_ERROR_INTERNAL_ERROR = 11, ///< Internal error
HIPRTC_ERROR_LINKING = 100 ///< Error in linking
} hiprtcResult;
/**
* hiprtc JIT option
*/
#define hiprtcJIT_option hipJitOption
#define HIPRTC_JIT_MAX_REGISTERS \
hipJitOptionMaxRegisters ///< CUDA Only Maximum registers may be used in a
///< thread, passed to compiler
#define HIPRTC_JIT_THREADS_PER_BLOCK \
hipJitOptionThreadsPerBlock ///< CUDA Only Number of thread per block
#define HIPRTC_JIT_WALL_TIME hipJitOptionWallTime ///< CUDA Only Value for total wall clock time
#define HIPRTC_JIT_INFO_LOG_BUFFER \
hipJitOptionInfoLogBuffer ///< CUDA Only Pointer to the buffer with
///< logged information
#define HIPRTC_JIT_INFO_LOG_BUFFER_SIZE_BYTES \
hipJitOptionInfoLogBufferSizeBytes ///< CUDA Only Size of the buffer
///< in bytes for logged info
#define HIPRTC_JIT_ERROR_LOG_BUFFER \
hipJitOptionErrorLogBuffer ///< CUDA Only Pointer to the buffer
///< with logged error(s)
#define HIPRTC_JIT_ERROR_LOG_BUFFER_SIZE_BYTES \
hipJitOptionErrorLogBufferSizeBytes ///< CUDA Only Size of the buffer in
///< bytes for logged error(s)
#define HIPRTC_JIT_OPTIMIZATION_LEVEL \
hipJitOptionOptimizationLevel ///< Value of optimization level for
///< generated codes, acceptable
///< options -O0, -O1, -O2, -O3
#define HIPRTC_JIT_TARGET_FROM_HIPCONTEXT \
hipJitOptionTargetFromContext ///< CUDA Only The target context,
///< which is the default
#define HIPRTC_JIT_TARGET hipJitOptionTarget ///< CUDA Only JIT target
#define HIPRTC_JIT_FALLBACK_STRATEGY hipJitOptionFallbackStrategy ///< CUDA Only Fallback strategy
#define HIPRTC_JIT_GENERATE_DEBUG_INFO \
hipJitOptionGenerateDebugInfo ///< CUDA Only Generate debug information
#define HIPRTC_JIT_LOG_VERBOSE hipJitOptionLogVerbose ///< CUDA Only Generate log verbose
#define HIPRTC_JIT_GENERATE_LINE_INFO \
hipJitOptionGenerateLineInfo ///< CUDA Only Generate line number information
#define HIPRTC_JIT_CACHE_MODE hipJitOptionCacheMode ///< CUDA Only Set cache mode
#define HIPRTC_JIT_NEW_SM3X_OPT hipJitOptionSm3xOpt ///< @deprecated CUDA Only New SM3X option.
#define HIPRTC_JIT_FAST_COMPILE hipJitOptionFastCompile ///< CUDA Only Set fast compile
#define HIPRTC_JIT_GLOBAL_SYMBOL_NAMES \
hipJitOptionGlobalSymbolNames ///< CUDA Only Array of device symbol names to be
///< relocated to the host
#define HIPRTC_JIT_GLOBAL_SYMBOL_ADDRESS \
hipJitOptionGlobalSymbolAddresses ///< CUDA Only Array of host addresses to be
///< relocated to the device
#define HIPRTC_JIT_GLOBAL_SYMBOL_COUNT \
hipJitOptionGlobalSymbolCount ///< CUDA Only Number of symbol count.
#define HIPRTC_JIT_LTO \
hipJitOptionLto ///< @deprecated CUDA Only Enable link-time
///< optimization for device code
#define HIPRTC_JIT_FTZ \
hipJitOptionFtz ///< @deprecated CUDA Only Set
///< single-precision denormals.
#define HIPRTC_JIT_PREC_DIV \
hipJitOptionPrecDiv ///< @deprecated CUDA Only Set
///< single-precision floating-point division
///< and reciprocals
#define HIPRTC_JIT_PREC_SQRT \
hipJitOptionPrecSqrt ///< @deprecated CUDA Only Set
///< single-precision floating-point
///< square root
#define HIPRTC_JIT_FMA \
hipJitOptionFma ///< @deprecated CUDA Only Enable
///< floating-point multiplies and
///< adds/subtracts operations
#define HIPRTC_JIT_POSITION_INDEPENDENT_CODE \
hipJitOptionPositionIndependentCode ///< CUDA Only Generates
///< Position Independent code
#define HIPRTC_JIT_MIN_CTA_PER_SM \
hipJitOptionMinCTAPerSM ///< CUDA Only Hints to JIT compiler
///< the minimum number of CTAs frin
///< kernel's grid to be mapped to SM
#define HIPRTC_JIT_MAX_THREADS_PER_BLOCK \
hipJitOptionMaxThreadsPerBlock ///< CUDA only Maximum number of
///< threads in a thread block
#define HIPRTC_JIT_OVERRIDE_DIRECT_VALUES \
hipJitOptionOverrideDirectiveValues ///< CUDA only Override Directive
///< Values
#define HIPRTC_JIT_NUM_OPTIONS hipJitOptionNumOptions ///< Number of options
#define HIPRTC_JIT_IR_TO_ISA_OPT_EXT \
hipJitOptionIRtoISAOptExt ///< HIP Only Linker options to be
///< passed on to compiler
#define HIPRTC_JIT_IR_TO_ISA_OPT_COUNT_EXT \
hipJitOptionIRtoISAOptCountExt ///< HIP Only Count of linker options
///< to be passed on to
/**
* hiprtc JIT input type
*/
#define hiprtcJITInputType hipJitInputType
#define HIPRTC_JIT_INPUT_CUBIN hipJitInputCubin ///< Cuda only Input Cubin
#define HIPRTC_JIT_INPUT_PTX hipJitInputPtx ///< Cuda only Input PTX
#define HIPRTC_JIT_INPUT_FATBINARY hipJitInputFatBinary ///< Cuda Only Input FAT Binary
#define HIPRTC_JIT_INPUT_OBJECT \
hipJitInputObject ///< Cuda Only Host Object with embedded device code
#define HIPRTC_JIT_INPUT_LIBRARY \
hipJitInputLibrary ///< Cuda Only Archive of Host Objects with embedded device code
#define HIPRTC_JIT_INPUT_NVVM \
hipJitInputNvvm ///< @deprecated CUDA only High Level intermediate code for LTO
#define HIPRTC_JIT_NUM_LEGACY_INPUT_TYPES \
hipJitNumLegacyInputTypes ///< Count of Legacy Input Types
#define HIPRTC_JIT_INPUT_LLVM_BITCODE \
hipJitInputLLVMBitcode ///< HIP Only LLVM Bitcode or IR assembly
#define HIPRTC_JIT_INPUT_LLVM_BUNDLED_BITCODE \
hipJitInputLLVMBundledBitcode ///< HIP Only LLVM Clang Bundled Code
#define HIPRTC_JIT_INPUT_LLVM_ARCHIVES_OF_BUNDLED_BITCODE \
hipJitInputLLVMArchivesOfBundledBitcode ///< HIP Only LLVM
///< Archives of
///< Bundled Bitcode
#define HIPRTC_JIT_INPUT_SPIRV hipJitInputSpirv ///< HIP Only SPIRV Code Object
#define HIPRTC_JIT_NUM_INPUT_TYPES hipJitNumInputTypes ///< Count of Input Types
/**
* @}
*/
/**
* hiprtc link state
*
*/
typedef struct ihiprtcLinkState* hiprtcLinkState;
/**
* @ingroup Runtime
*
* @brief Returns text string message to explain the error which occurred
*
* @param [in] result code to convert to string.
* @returns const char pointer to the NULL-terminated error string
*
* @warning In HIP, this function returns the name of the error,
* if the hiprtc result is defined, it will return "Invalid HIPRTC error code"
*
* @see hiprtcResult
*/
const char* hiprtcGetErrorString(hiprtcResult result);
/**
* @ingroup Runtime
* @brief Sets the parameters as major and minor version.
*
* @param [out] major HIP Runtime Compilation major version.
* @param [out] minor HIP Runtime Compilation minor version.
*
* @returns #HIPRTC_ERROR_INVALID_INPUT, #HIPRTC_SUCCESS
*
*/
hiprtcResult hiprtcVersion(int* major, int* minor);
/**
* hiprtc program
*
*/
typedef struct _hiprtcProgram* hiprtcProgram;
/**
* @ingroup Runtime
* @brief Adds the given name exprssion to the runtime compilation program.
*
* @param [in] prog runtime compilation program instance.
* @param [in] name_expression const char pointer to the name expression.
* @returns #HIPRTC_SUCCESS
*
* If const char pointer is NULL, it will return #HIPRTC_ERROR_INVALID_INPUT.
*
* @see hiprtcResult
*/
hiprtcResult hiprtcAddNameExpression(hiprtcProgram prog, const char* name_expression);
/**
* @ingroup Runtime
* @brief Compiles the given runtime compilation program.
*
* @param [in] prog runtime compilation program instance.
* @param [in] numOptions number of compiler options.
* @param [in] options compiler options as const array of strins.
* @returns #HIPRTC_SUCCESS
*
* If the compiler failed to build the runtime compilation program,
* it will return #HIPRTC_ERROR_COMPILATION.
*
* @see hiprtcResult
*/
hiprtcResult hiprtcCompileProgram(hiprtcProgram prog, int numOptions, const char* const* options);
/**
* @ingroup Runtime
* @brief Creates an instance of hiprtcProgram with the given input parameters,
* and sets the output hiprtcProgram prog with it.
*
* @param [in, out] prog runtime compilation program instance.
* @param [in] src const char pointer to the program source.
* @param [in] name const char pointer to the program name.
* @param [in] numHeaders number of headers.
* @param [in] headers array of strings pointing to headers.
* @param [in] includeNames array of strings pointing to names included in program source.
* @returns #HIPRTC_SUCCESS
*
* Any invalide input parameter, it will return #HIPRTC_ERROR_INVALID_INPUT
* or #HIPRTC_ERROR_INVALID_PROGRAM.
*
* If failed to create the program, it will return #HIPRTC_ERROR_PROGRAM_CREATION_FAILURE.
*
* @see hiprtcResult
*/
hiprtcResult hiprtcCreateProgram(hiprtcProgram* prog, const char* src, const char* name,
int numHeaders, const char* const* headers,
const char* const* includeNames);
/**
* @brief Destroys an instance of given hiprtcProgram.
* @ingroup Runtime
* @param [in] prog runtime compilation program instance.
* @returns #HIPRTC_SUCCESS
*
* If prog is NULL, it will return #HIPRTC_ERROR_INVALID_INPUT.
*
* @see hiprtcResult
*/
hiprtcResult hiprtcDestroyProgram(hiprtcProgram* prog);
/**
* @brief Gets the lowered (mangled) name from an instance of hiprtcProgram with the given input
* parameters, and sets the output lowered_name with it.
* @ingroup Runtime
* @param [in] prog runtime compilation program instance.
* @param [in] name_expression const char pointer to the name expression.
* @param [in, out] lowered_name const char array to the lowered (mangled) name.
* @returns #HIPRTC_SUCCESS
*
* If any invalide nullptr input parameters, it will return #HIPRTC_ERROR_INVALID_INPUT
*
* If name_expression is not found, it will return #HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID
*
* If failed to get lowered_name from the program, it will return #HIPRTC_ERROR_COMPILATION.
*
* @see hiprtcResult
*/
hiprtcResult hiprtcGetLoweredName(hiprtcProgram prog, const char* name_expression,
const char** lowered_name);
/**
* @brief Gets the log generated by the runtime compilation program instance.
* @ingroup Runtime
* @param [in] prog runtime compilation program instance.
* @param [out] log memory pointer to the generated log.
* @returns #HIPRTC_SUCCESS
*
* @see hiprtcResult
*/
hiprtcResult hiprtcGetProgramLog(hiprtcProgram prog, char* log);
/**
* @brief Gets the size of log generated by the runtime compilation program instance.
*
* @param [in] prog runtime compilation program instance.
* @param [out] logSizeRet size of generated log.
* @returns #HIPRTC_SUCCESS
*
* @see hiprtcResult
*/
hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram prog, size_t* logSizeRet);
/**
* @brief Gets the pointer of compilation binary by the runtime compilation program instance.
* @ingroup Runtime
* @param [in] prog runtime compilation program instance.
* @param [out] code char pointer to binary.
* @returns #HIPRTC_SUCCESS
*
* @see hiprtcResult
*/
hiprtcResult hiprtcGetCode(hiprtcProgram prog, char* code);
/**
* @brief Gets the size of compilation binary by the runtime compilation program instance.
* @ingroup Runtime
* @param [in] prog runtime compilation program instance.
* @param [out] codeSizeRet the size of binary.
* @returns #HIPRTC_SUCCESS
*
* @see hiprtcResult
*/
hiprtcResult hiprtcGetCodeSize(hiprtcProgram prog, size_t* codeSizeRet);
/**
* @brief Gets the pointer of compiled bitcode by the runtime compilation program instance.
*
* @param [in] prog runtime compilation program instance.
* @param [out] bitcode char pointer to bitcode.
* @return HIPRTC_SUCCESS
*
* @see hiprtcResult
*/
hiprtcResult hiprtcGetBitcode(hiprtcProgram prog, char* bitcode);
/**
* @brief Gets the size of compiled bitcode by the runtime compilation program instance.
* @ingroup Runtime
*
* @param [in] prog runtime compilation program instance.
* @param [out] bitcode_size the size of bitcode.
* @returns #HIPRTC_SUCCESS
*
* @see hiprtcResult
*/
hiprtcResult hiprtcGetBitcodeSize(hiprtcProgram prog, size_t* bitcode_size);
/**
* @brief Creates the link instance via hiprtc APIs.
* @ingroup Runtime
* @param [in] num_options Number of options
* @param [in] option_ptr Array of options
* @param [in] option_vals_pptr Array of option values cast to void*
* @param [out] hip_link_state_ptr hiprtc link state created upon success
*
* @returns #HIPRTC_SUCCESS, #HIPRTC_ERROR_INVALID_INPUT, #HIPRTC_ERROR_INVALID_OPTION
*
* @see hiprtcResult
*/
hiprtcResult hiprtcLinkCreate(unsigned int num_options, hiprtcJIT_option* option_ptr,
void** option_vals_pptr, hiprtcLinkState* hip_link_state_ptr);
/**
* @brief Adds a file with bit code to be linked with options
* @ingroup Runtime
* @param [in] hip_link_state hiprtc link state
* @param [in] input_type Type of the input data or bitcode
* @param [in] file_path Path to the input file where bitcode is present
* @param [in] num_options Size of the options
* @param [in] options_ptr Array of options applied to this input
* @param [in] option_values Array of option values cast to void*
*
* @returns #HIPRTC_SUCCESS
*
* If input values are invalid, it will
* @return #HIPRTC_ERROR_INVALID_INPUT
*
* @see hiprtcResult
*/
hiprtcResult hiprtcLinkAddFile(hiprtcLinkState hip_link_state, hiprtcJITInputType input_type,
const char* file_path, unsigned int num_options,
hiprtcJIT_option* options_ptr, void** option_values);
/**
* @brief Completes the linking of the given program.
* @ingroup Runtime
* @param [in] hip_link_state hiprtc link state
* @param [in] input_type Type of the input data or bitcode
* @param [in] image Input data which is null terminated
* @param [in] image_size Size of the input data
* @param [in] name Optional name for this input
* @param [in] num_options Size of the options
* @param [in] options_ptr Array of options applied to this input
* @param [in] option_values Array of option values cast to void*
*
* @returns #HIPRTC_SUCCESS, #HIPRTC_ERROR_INVALID_INPUT
*
* If adding the file fails, it will
* @return #HIPRTC_ERROR_PROGRAM_CREATION_FAILURE
*
* @see hiprtcResult
*/
hiprtcResult hiprtcLinkAddData(hiprtcLinkState hip_link_state, hiprtcJITInputType input_type,
void* image, size_t image_size, const char* name,
unsigned int num_options, hiprtcJIT_option* options_ptr,
void** option_values);
/**
* @brief Completes the linking of the given program.
* @ingroup Runtime
* @param [in] hip_link_state hiprtc link state
* @param [out] bin_out Upon success, points to the output binary
* @param [out] size_out Size of the binary is stored (optional)
*
* @returns #HIPRTC_SUCCESS
*
* If adding the data fails, it will
* @return #HIPRTC_ERROR_LINKING
*
* @see hiprtcResult
*/
hiprtcResult hiprtcLinkComplete(hiprtcLinkState hip_link_state, void** bin_out, size_t* size_out);
/**
* @brief Deletes the link instance via hiprtc APIs.
* @ingroup Runtime
* @param [in] hip_link_state link state instance
*
* @returns #HIPRTC_SUCCESS
*
* @see hiprtcResult
*/
hiprtcResult hiprtcLinkDestroy(hiprtcLinkState hip_link_state);
#if !defined(_WIN32)
#pragma GCC visibility pop
#endif
#ifdef __cplusplus
}
#endif /* __cplusplus */
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif

84
external/hip/library_types.h vendored Normal file
View File

@@ -0,0 +1,84 @@
/*
Copyright (c) 2015 - 2023 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 HIP_INCLUDE_HIP_LIBRARY_TYPES_H
#define HIP_INCLUDE_HIP_LIBRARY_TYPES_H
#if !defined(__HIPCC_RTC__)
#include <hip/hip_common.h>
#endif
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
typedef enum hipDataType {
HIP_R_32F = 0,
HIP_R_64F = 1,
HIP_R_16F = 2,
HIP_R_8I = 3,
HIP_C_32F = 4,
HIP_C_64F = 5,
HIP_C_16F = 6,
HIP_C_8I = 7,
HIP_R_8U = 8,
HIP_C_8U = 9,
HIP_R_32I = 10,
HIP_C_32I = 11,
HIP_R_32U = 12,
HIP_C_32U = 13,
HIP_R_16BF = 14,
HIP_C_16BF = 15,
HIP_R_4I = 16,
HIP_C_4I = 17,
HIP_R_4U = 18,
HIP_C_4U = 19,
HIP_R_16I = 20,
HIP_C_16I = 21,
HIP_R_16U = 22,
HIP_C_16U = 23,
HIP_R_64I = 24,
HIP_C_64I = 25,
HIP_R_64U = 26,
HIP_C_64U = 27,
HIP_R_8F_E4M3 = 28,
HIP_R_8F_E5M2 = 29,
HIP_R_8F_UE8M0 = 30,
HIP_R_6F_E2M3 = 31,
HIP_R_6F_E3M2 = 32,
HIP_R_4F_E2M1 = 33,
// HIP specific Data Types
HIP_R_8F_E4M3_FNUZ = 1000,
HIP_R_8F_E5M2_FNUZ = 1001,
} hipDataType;
typedef enum hipLibraryPropertyType {
HIP_LIBRARY_MAJOR_VERSION,
HIP_LIBRARY_MINOR_VERSION,
HIP_LIBRARY_PATCH_LEVEL
} hipLibraryPropertyType;
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include "library_types.h"
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif

138
external/hip/linker_types.h vendored Executable file
View File

@@ -0,0 +1,138 @@
/*
Copyright (c) 2015 - 2023 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 HIP_INCLUDE_HIP_LINKER_TYPES_H
#define HIP_INCLUDE_HIP_LINKER_TYPES_H
#if defined(__clang__)
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wreserved-identifier"
#pragma clang diagnostic ignored "-Wreserved-macro-identifier"
#endif
#if !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#elif defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
/**
* @defgroup LinkerTypes Jit Linker Data Types
* @{
* This section describes the Jit Linker data types.
*
*/
/**
* hipJitOption
*/
typedef enum hipJitOption {
hipJitOptionMaxRegisters = 0, ///< CUDA Only Maximum registers may be used in a thread,
///< passed to compiler
hipJitOptionThreadsPerBlock, ///< CUDA Only Number of thread per block
hipJitOptionWallTime, ///< CUDA Only Value for total wall clock time
hipJitOptionInfoLogBuffer, ///< CUDA Only Pointer to the buffer with logged information
hipJitOptionInfoLogBufferSizeBytes, ///< CUDA Only Size of the buffer in bytes for logged info
hipJitOptionErrorLogBuffer, ///< CUDA Only Pointer to the buffer with logged error(s)
hipJitOptionErrorLogBufferSizeBytes, ///< CUDA Only Size of the buffer in bytes for logged
///< error(s)
hipJitOptionOptimizationLevel, ///< Value of optimization level for generated codes, acceptable
///< options -O0, -O1, -O2, -O3
hipJitOptionTargetFromContext, ///< CUDA Only The target context, which is the default
hipJitOptionTarget, ///< CUDA Only JIT target
hipJitOptionFallbackStrategy, ///< CUDA Only Fallback strategy
hipJitOptionGenerateDebugInfo, ///< CUDA Only Generate debug information
hipJitOptionLogVerbose, ///< CUDA Only Generate log verbose
hipJitOptionGenerateLineInfo, ///< CUDA Only Generate line number information
hipJitOptionCacheMode, ///< CUDA Only Set cache mode
hipJitOptionSm3xOpt, ///< @deprecated CUDA Only New SM3X option.
hipJitOptionFastCompile, ///< CUDA Only Set fast compile
hipJitOptionGlobalSymbolNames, ///< CUDA Only Array of device symbol names to be relocated to the
///< host
hipJitOptionGlobalSymbolAddresses, ///< CUDA Only Array of host addresses to be relocated to the
///< device
hipJitOptionGlobalSymbolCount, ///< CUDA Only Number of symbol count.
hipJitOptionLto, ///< @deprecated CUDA Only Enable link-time optimization for device code
hipJitOptionFtz, ///< @deprecated CUDA Only Set single-precision denormals.
hipJitOptionPrecDiv, ///< @deprecated CUDA Only Set single-precision floating-point division
///< and reciprocals
hipJitOptionPrecSqrt, ///< @deprecated CUDA Only Set single-precision floating-point square root
hipJitOptionFma, ///< @deprecated CUDA Only Enable floating-point multiplies and
///< adds/subtracts operations
hipJitOptionPositionIndependentCode, ///< CUDA Only Generates Position Independent code
hipJitOptionMinCTAPerSM, ///< CUDA Only Hints to JIT compiler the minimum number of CTAs frin
///< kernel's grid to be mapped to SM
hipJitOptionMaxThreadsPerBlock, ///< CUDA only Maximum number of threads in a thread block
hipJitOptionOverrideDirectiveValues, ///< Cuda only Override Directive values
hipJitOptionNumOptions, ///< Number of options
hipJitOptionIRtoISAOptExt = 10000, ///< Hip Only Linker options to be passed on to compiler
hipJitOptionIRtoISAOptCountExt, ///< Hip Only Count of linker options to be passed on to compiler
} hipJitOption;
/**
* hipJitInputType
*/
typedef enum hipJitInputType {
hipJitInputCubin = 0, ///< Cuda only Input cubin
hipJitInputPtx, ///< Cuda only Input PTX
hipJitInputFatBinary, ///< Cuda Only Input FAT Binary
hipJitInputObject, ///< Cuda Only Host Object with embedded device code
hipJitInputLibrary, ///< Cuda Only Archive of Host Objects with embedded
///< device code
hipJitInputNvvm, ///< @deprecated Cuda only High Level intermediate
///< code for LTO
hipJitNumLegacyInputTypes, ///< Count of Legacy Input Types
hipJitInputLLVMBitcode = 100, ///< HIP Only LLVM Bitcode or IR assembly
hipJitInputLLVMBundledBitcode = 101, ///< HIP Only LLVM Clang Bundled Code
hipJitInputLLVMArchivesOfBundledBitcode = 102, ///< HIP Only LLVM Archive of Bundled Bitcode
hipJitInputSpirv = 103, ///< HIP Only SPIRV Code Object
hipJitNumInputTypes = 10 ///< Count of Input Types
} hipJitInputType;
/**
* hipJitCacheMode
*/
typedef enum hipJitCacheMode {
hipJitCacheOptionNone = 0,
hipJitCacheOptionCG,
hipJitCacheOptionCA
} hipJitCacheMode;
/**
* hipJitFallback
*/
typedef enum hipJitFallback {
hipJitPreferPTX = 0,
hipJitPreferBinary,
} hipJitFallback;
typedef enum hipLibraryOption_e {
hipLibraryHostUniversalFunctionAndDataTable = 0,
hipLibraryBinaryIsPreserved = 1
} hipLibraryOption;
// doxygen end LinkerTypes
/**
* @}
*/
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif // HIP_INCLUDE_HIP_LINKER_TYPES_H

42
external/hip/math_functions.h vendored Normal file
View File

@@ -0,0 +1,42 @@
/*
Copyright (c) 2015 - 2023 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 HIP_INCLUDE_HIP_MATH_FUNCTIONS_H
#define HIP_INCLUDE_HIP_MATH_FUNCTIONS_H
// Some standard header files, these are included by hc.hpp and so want to make them avail on both
// paths to provide a consistent include env and avoid "missing symbol" errors that only appears
// on NVCC path:
#if !defined(__HIPCC_RTC__)
#include <hip/hip_common.h>
#endif
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
#include <hip/amd_detail/amd_math_functions.h>
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
// #include <hip/nvidia_detail/math_functions.h>
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#endif

65
external/hip/surface_types.h vendored Normal file
View File

@@ -0,0 +1,65 @@
/*
Copyright (c) 2022 - 2023 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.
*/
/**
* @file surface_types.h
* @brief Defines surface types for HIP runtime.
*/
#ifndef HIP_INCLUDE_HIP_SURFACE_TYPES_H
#define HIP_INCLUDE_HIP_SURFACE_TYPES_H
#if defined(__clang__)
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wreserved-identifier"
#endif
#if !defined(__HIPCC_RTC__)
#include <hip/driver_types.h>
#endif
/**
* An opaque value that represents a hip surface object
*/
struct __hip_surface;
typedef struct __hip_surface* hipSurfaceObject_t;
/**
* hip surface reference
*/
struct surfaceReference {
hipSurfaceObject_t surfaceObject;
};
/**
* hip surface boundary modes
*/
enum hipSurfaceBoundaryMode {
hipBoundaryModeZero = 0,
hipBoundaryModeTrap = 1,
hipBoundaryModeClamp = 2
};
#if defined(__clang__)
#pragma clang diagnostic pop
#endif
#endif /* !HIP_INCLUDE_HIP_SURFACE_TYPES_H */

193
external/hip/texture_types.h vendored Normal file
View File

@@ -0,0 +1,193 @@
/*
Copyright (c) 2015 - 2023 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 HIP_INCLUDE_HIP_TEXTURE_TYPES_H
#define HIP_INCLUDE_HIP_TEXTURE_TYPES_H
#if defined(__clang__)
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wreserved-identifier"
#pragma clang diagnostic ignored "-Wreserved-macro-identifier"
#pragma clang diagnostic ignored "-Wc++98-compat"
#endif
#if !defined(__HIPCC_RTC__)
#include <hip/hip_common.h>
#endif
#if !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include "texture_types.h"
#elif defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
/*******************************************************************************
* *
* *
* *
*******************************************************************************/
#if !defined(__HIPCC_RTC__)
#include <hip/channel_descriptor.h>
#include <hip/driver_types.h>
#endif // !defined(__HIPCC_RTC__)
#define hipTextureType1D 0x01
#define hipTextureType2D 0x02
#define hipTextureType3D 0x03
#define hipTextureTypeCubemap 0x0C
#define hipTextureType1DLayered 0xF1
#define hipTextureType2DLayered 0xF2
#define hipTextureTypeCubemapLayered 0xFC
/**
* Should be same as HSA_IMAGE_OBJECT_SIZE_DWORD/HSA_SAMPLER_OBJECT_SIZE_DWORD
*/
#define HIP_IMAGE_OBJECT_SIZE_DWORD 12
#define HIP_SAMPLER_OBJECT_SIZE_DWORD 8
#define HIP_SAMPLER_OBJECT_OFFSET_DWORD HIP_IMAGE_OBJECT_SIZE_DWORD
#define HIP_TEXTURE_OBJECT_SIZE_DWORD (HIP_IMAGE_OBJECT_SIZE_DWORD + HIP_SAMPLER_OBJECT_SIZE_DWORD)
/**
* An opaque value that represents a hip texture object
*/
struct __hip_texture;
typedef struct __hip_texture* hipTextureObject_t;
/**
* hip texture address modes
*/
enum hipTextureAddressMode {
hipAddressModeWrap = 0,
hipAddressModeClamp = 1,
hipAddressModeMirror = 2,
hipAddressModeBorder = 3
};
/**
* hip texture filter modes
*/
enum hipTextureFilterMode { hipFilterModePoint = 0, hipFilterModeLinear = 1 };
/**
* hip texture read modes
*/
enum hipTextureReadMode { hipReadModeElementType = 0, hipReadModeNormalizedFloat = 1 };
/**
* hip texture reference
*/
typedef struct textureReference {
int normalized;
enum hipTextureReadMode readMode; // used only for driver API's
enum hipTextureFilterMode filterMode;
enum hipTextureAddressMode addressMode[3]; // Texture address mode for up to 3 dimensions
struct hipChannelFormatDesc channelDesc;
int sRGB; // Perform sRGB->linear conversion during texture read
unsigned int maxAnisotropy; // Limit to the anisotropy ratio
enum hipTextureFilterMode mipmapFilterMode;
float mipmapLevelBias;
float minMipmapLevelClamp;
float maxMipmapLevelClamp;
hipTextureObject_t textureObject;
int numChannels;
enum hipArray_Format format;
} textureReference;
/**
* hip texture descriptor
*/
typedef struct hipTextureDesc {
enum hipTextureAddressMode addressMode[3]; // Texture address mode for up to 3 dimensions
enum hipTextureFilterMode filterMode;
enum hipTextureReadMode readMode;
int sRGB; // Perform sRGB->linear conversion during texture read
float borderColor[4];
int normalizedCoords;
unsigned int maxAnisotropy;
enum hipTextureFilterMode mipmapFilterMode;
float mipmapLevelBias;
float minMipmapLevelClamp;
float maxMipmapLevelClamp;
} hipTextureDesc;
#if __cplusplus
/*******************************************************************************
* *
* *
* *
*******************************************************************************/
#if __HIP__
#define __HIP_TEXTURE_ATTRIB __attribute__((device_builtin_texture_type))
#else
#define __HIP_TEXTURE_ATTRIB
#endif
typedef textureReference* hipTexRef;
template <class T, int texType = hipTextureType1D,
enum hipTextureReadMode mode = hipReadModeElementType>
struct __HIP_TEXTURE_ATTRIB texture : public textureReference {
texture(int norm = 0, enum hipTextureFilterMode fMode = hipFilterModePoint,
enum hipTextureAddressMode aMode = hipAddressModeClamp) {
normalized = norm;
readMode = mode;
filterMode = fMode;
addressMode[0] = aMode;
addressMode[1] = aMode;
addressMode[2] = aMode;
channelDesc = hipCreateChannelDesc<T>();
sRGB = 0;
textureObject = nullptr;
maxAnisotropy = 0;
mipmapLevelBias = 0;
minMipmapLevelClamp = 0;
maxMipmapLevelClamp = 0;
}
texture(int norm, enum hipTextureFilterMode fMode, enum hipTextureAddressMode aMode,
struct hipChannelFormatDesc desc) {
normalized = norm;
readMode = mode;
filterMode = fMode;
addressMode[0] = aMode;
addressMode[1] = aMode;
addressMode[2] = aMode;
channelDesc = desc;
sRGB = 0;
textureObject = nullptr;
maxAnisotropy = 0;
mipmapLevelBias = 0;
minMipmapLevelClamp = 0;
maxMipmapLevelClamp = 0;
}
};
#endif /* __cplusplus */
#else
#error ("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
#if defined(__clang__)
#pragma clang diagnostic pop
#endif
#endif

1
external/spdlog vendored Submodule

Submodule external/spdlog added at 472945ba48

10
kernels/kernels.hip Normal file
View File

@@ -0,0 +1,10 @@
#include <hip/hip_runtime.h>
extern "C" __global__
void saxpy(float* y, const float* x, float a, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
y[i] += a * x[i];
}
}

BIN
kernels/kernels_gfx1100.co Normal file
View File

Binary file not shown.

82
src/main.cpp Normal file
View File

@@ -0,0 +1,82 @@
#include <hip/hip_runtime.h>
#include <iostream>
#include <vector>
int main() {
int N = 1 << 20;
// -----------------------------
// Host data
// -----------------------------
std::vector<float> hx(N, 1.0f);
std::vector<float> hy(N, 2.0f);
// -----------------------------
// Device allocation
// -----------------------------
float* dx = nullptr;
float* dy = nullptr;
hipInit(0);
hipMalloc(&dx, N * sizeof(float));
hipMalloc(&dy, N * sizeof(float));
hipMemcpy(dx, hx.data(), N * sizeof(float), hipMemcpyHostToDevice);
hipMemcpy(dy, hy.data(), N * sizeof(float), hipMemcpyHostToDevice);
// -----------------------------
// Load precompiled kernel
// -----------------------------
hipModule_t module{};
hipFunction_t kernel{};
hipModuleLoad(&module, "kernels/kernels_gfx1100.co");
hipModuleGetFunction(&kernel, module, "saxpy");
// -----------------------------
// Kernel args
// -----------------------------
float a = 3.0f;
void* args[] = {
&dy,
&dx,
&a,
&N
};
// -----------------------------
// Launch
// -----------------------------
const int blockSize = 256;
const int gridSize = (N + blockSize - 1) / blockSize;
hipModuleLaunchKernel(
kernel,
gridSize, 1, 1,
blockSize, 1, 1,
0,
nullptr,
args,
nullptr
);
hipDeviceSynchronize();
// -----------------------------
// Copy back
// -----------------------------
hipMemcpy(hy.data(), dy, N * sizeof(float), hipMemcpyDeviceToHost);
std::cout << "hy[0] = " << hy[0] << std::endl;
// -----------------------------
// Cleanup
// -----------------------------
hipFree(dx);
hipFree(dy);
hipModuleUnload(module);
return 0;
}