develop/amd_openvx/openvx/include/vx_ext_amd.h Source File

develop/amd_openvx/openvx/include/vx_ext_amd.h Source File#

MIVisionX: develop/amd_openvx/openvx/include/vx_ext_amd.h Source File
 /*
 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 _VX_EXT_AMD_H_
 #define _VX_EXT_AMD_H_
  
 #include <VX/vx.h>
 #include <VX/vx_khr_nn.h>
 #ifdef __cplusplus
 #include <string>
 #endif
  
 #define AGO_TARGET_AFFINITY_CPU 0x0010 // CPU
 #define AGO_TARGET_AFFINITY_GPU 0x0020 // GPU
 #define AGO_TARGET_AFFINITY_APU 0x0030 // APU
  
 #define AGO_MAX_PARAMS 32
 #define AGO_MERGE_RULE_MAX_FIND 4
 #define AGO_MERGE_RULE_MAX_REPLACE 4
 #define AGO_MERGE_RULE_SOLITARY_FLAG 0x20
 #define AGO_TARGET_AFFINITY_GPU_INFO_DEVICE_MASK 0x0F
 #define AGO_TARGET_AFFINITY_GPU_INFO_SVM_MASK 0xF0
 #define AGO_TARGET_AFFINITY_GPU_INFO_SVM_ENABLE 0x10
 #define AGO_TARGET_AFFINITY_GPU_INFO_SVM_AS_CLMEM 0x20
 #define AGO_TARGET_AFFINITY_GPU_INFO_SVM_NO_FGS 0x40
  
 #define VX_MAX_STRING_BUFFER_SIZE_AMD 256
  
 #define VX_NN_ACTIVATION_LEAKY_RELU (VX_ENUM_BASE(VX_ID_AMD, VX_ENUM_NN_ACTIVATION_FUNCTION_TYPE) + 0x9)
  
 #define VX_CONTEXT_ATTRIBUTE_NONLINEAR_MAX_DIMENSION VX_CONTEXT_NONLINEAR_MAX_DIMENSION
 #define VX_CONTEXT_ATTRIBUTE_IMMEDIATE_BORDER_POLICY VX_CONTEXT_IMMEDIATE_BORDER_POLICY
  
 #define VX_GRAPH_ATTRIBUTE_STATE VX_GRAPH_STATE
  
 #define VX_OBJECT_ARRAY_ATTRIBUTE_ITEMTYPE VX_OBJECT_ARRAY_ITEMTYPE
 #define VX_OBJECT_ARRAY_ATTRIBUTE_NUMITEMS VX_OBJECT_ARRAY_NUMITEMS
  
 enum ago_type_public_e
 {
     VX_TYPE_STRING_AMD = 0x011,
     AGO_TYPE_KEYPOINT_XYS = VX_TYPE_VENDOR_STRUCT_START,
     AGO_TYPE_MEANSTDDEV_DATA = VX_TYPE_VENDOR_OBJECT_START,
     AGO_TYPE_MINMAXLOC_DATA,
     AGO_TYPE_CANNY_STACK,
     AGO_TYPE_SCALE_MATRIX,
 };
  
 enum vx_context_attribute_amd_e
 {
     VX_CONTEXT_ATTRIBUTE_AMD_OPENCL_CONTEXT = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_CONTEXT) + 0x01,
     VX_CONTEXT_ATTRIBUTE_AMD_AFFINITY = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_CONTEXT) + 0x02,
     VX_CONTEXT_ATTRIBUTE_AMD_SET_TEXT_MACRO = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_CONTEXT) + 0x03,
     VX_CONTEXT_ATTRIBUTE_AMD_SET_MERGE_RULE = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_CONTEXT) + 0x04,
     VX_CONTEXT_MAX_TENSOR_DIMENSIONS = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_CONTEXT) + 0x05,
     VX_CONTEXT_CL_QUEUE_PROPERTIES = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_CONTEXT) + 0x06,
     VX_CONTEXT_ATTRIBUTE_AMD_HIP_DEVICE = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_CONTEXT) + 0x07,
 };
  
 enum vx_kernel_attribute_amd_e
 {
     VX_KERNEL_ATTRIBUTE_AMD_QUERY_TARGET_SUPPORT = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_KERNEL) + 0x01,
     VX_KERNEL_ATTRIBUTE_AMD_OPENCL_CODEGEN_CALLBACK = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_KERNEL) + 0x02,
     VX_KERNEL_ATTRIBUTE_AMD_NODE_REGEN_CALLBACK = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_KERNEL) + 0x03,
     VX_KERNEL_ATTRIBUTE_AMD_OPENCL_GLOBAL_WORK_UPDATE_CALLBACK = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_KERNEL) + 0x04,
     VX_KERNEL_ATTRIBUTE_AMD_GPU_BUFFER_ACCESS_ENABLE = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_KERNEL) + 0x05,
     VX_KERNEL_ATTRIBUTE_AMD_GPU_BUFFER_UPDATE_CALLBACK = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_KERNEL) + 0x06,
 };
  
 enum vx_graph_attribute_amd_e
 {
     VX_GRAPH_ATTRIBUTE_AMD_AFFINITY = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_GRAPH) + 0x01,
     VX_GRAPH_ATTRIBUTE_AMD_IMPORT_FROM_TEXT = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_GRAPH) + 0x02,
     VX_GRAPH_ATTRIBUTE_AMD_EXPORT_TO_TEXT = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_GRAPH) + 0x03,
     VX_GRAPH_ATTRIBUTE_AMD_OPTIMIZER_FLAGS = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_GRAPH) + 0x04,
     VX_GRAPH_ATTRIBUTE_AMD_PERFORMANCE_INTERNAL_LAST = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_GRAPH) + 0x05,
     VX_GRAPH_ATTRIBUTE_AMD_PERFORMANCE_INTERNAL_AVG = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_GRAPH) + 0x06,
     VX_GRAPH_ATTRIBUTE_AMD_PERFORMANCE_INTERNAL_PROFILE = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_GRAPH) + 0x07,
     VX_GRAPH_ATTRIBUTE_AMD_OPENCL_COMMAND_QUEUE = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_GRAPH) + 0x08,
     VX_GRAPH_ATTRIBUTE_AMD_CPU_NUM_THREADS = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_GRAPH) + 0x09
 };
  
 enum vx_node_attribute_amd_e
 {
     VX_NODE_ATTRIBUTE_AMD_AFFINITY = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_NODE) + 0x01,
     VX_NODE_ATTRIBUTE_AMD_OPENCL_COMMAND_QUEUE = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_NODE) + 0x02,
     VX_NODE_ATTRIBUTE_AMD_HIP_STREAM = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_NODE) + 0x03,
     VX_NODE_ATTRIBUTE_AMD_CPU_NUM_THREADS = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_NODE) + 0x04
 };
  
 enum vx_image_attribute_amd_e
 {
     VX_IMAGE_ATTRIBUTE_AMD_OPENCL_BUFFER = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_IMAGE) + 0x01,
     VX_IMAGE_ATTRIBUTE_AMD_GPU_BUFFER_OFFSET = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_IMAGE) + 0x02,
     VX_IMAGE_ATTRIBUTE_AMD_ENABLE_USER_BUFFER_GPU = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_IMAGE) + 0x03,
     VX_IMAGE_ATTRIBUTE_AMD_GPU_BUFFER_STRIDE = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_IMAGE) + 0x04,
     VX_IMAGE_ATTRIBUTE_AMD_HOST_BUFFER = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_IMAGE) + 0x05,
     VX_IMAGE_ATTRIBUTE_AMD_HIP_BUFFER = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_IMAGE) + 0x06,
 };
  
 enum vx_tensor_attribute_amd_e
 {
     VX_TENSOR_STRIDE_GPU = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_TENSOR) + 0x5,
     VX_TENSOR_OFFSET_GPU = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_TENSOR) + 0x6,
     VX_TENSOR_BUFFER_OPENCL = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_TENSOR) + 0x7,
     VX_TENSOR_BUFFER_HIP = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_TENSOR) + 0x8,
     VX_TENSOR_MEMORY_TYPE = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_TENSOR) + 0x9,
     VX_TENSOR_BUFFER_HOST = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_TENSOR) + 0x10,
 };
  
 enum vx_array_attribute_amd_e
 {
     VX_ARRAY_BUFFER_OPENCL = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_ARRAY) + 0x9,
     VX_ARRAY_BUFFER_HIP = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_ARRAY) + 0x10,
     VX_ARRAY_BUFFER = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_ARRAY) + 0x11,
     VX_ARRAY_OFFSET_GPU = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_ARRAY) + 0x12
 };
  
 enum vx_directive_amd_e
 {
     VX_DIRECTIVE_AMD_READ_ONLY = VX_ENUM_BASE(VX_ID_AMD, VX_ENUM_DIRECTIVE) + 0x01,
     VX_DIRECTIVE_AMD_COPY_TO_OPENCL = VX_ENUM_BASE(VX_ID_AMD, VX_ENUM_DIRECTIVE) + 0x02,
     VX_DIRECTIVE_AMD_ENABLE_PROFILE_CAPTURE = VX_ENUM_BASE(VX_ID_AMD, VX_ENUM_DIRECTIVE) + 0x03,
     VX_DIRECTIVE_AMD_DISABLE_PROFILE_CAPTURE = VX_ENUM_BASE(VX_ID_AMD, VX_ENUM_DIRECTIVE) + 0x04,
     VX_DIRECTIVE_AMD_DISABLE_GPU_FLUSH = VX_ENUM_BASE(VX_ID_AMD, VX_ENUM_DIRECTIVE) + 0x05,
     VX_DIRECTIVE_AMD_COPY_TO_HIPMEM = VX_ENUM_BASE(VX_ID_AMD, VX_ENUM_DIRECTIVE) + 0x02,
 };
  
 enum vx_memory_type_amd_e
 {
     VX_MEMORY_TYPE_OPENCL = VX_ENUM_BASE(VX_ID_KHRONOS, VX_ENUM_MEMORY_TYPE) + 0x2,
     VX_MEMORY_TYPE_HIP = VX_ENUM_BASE(VX_ID_KHRONOS, VX_ENUM_MEMORY_TYPE) + 0x3,
 };
  
 enum vx_color_space_amd_e
 {
     VX_COLOR_SPACE_BT2020 = VX_ENUM_BASE(VX_ID_AMD, VX_ENUM_COLOR_SPACE) + 0x1,
 };
  
 enum vx_df_image_amd_e
 {
     VX_DF_IMAGE_U1_AMD = VX_DF_IMAGE('U', '0', '0', '1'),
     VX_DF_IMAGE_F16_AMD = VX_DF_IMAGE('F', '0', '1', '6'),
     VX_DF_IMAGE_F32_AMD = VX_DF_IMAGE('F', '0', '3', '2'),
     VX_DF_IMAGE_F64_AMD = VX_DF_IMAGE('F', '0', '6', '4'),
     VX_DF_IMAGE_F32x3_AMD = VX_DF_IMAGE('F', '3', '3', '2'),
 };
  
 enum vx_scalar_attribute_amd_e
 {
     VX_SCALAR_BUFFER = VX_ATTRIBUTE_BASE(VX_ID_AMD, VX_TYPE_SCALAR) + 0x1,
 };
  
 typedef struct _vx_tensor_t *vx_tensor;
  
 typedef struct
 {
     vx_size components;
     vx_size planes;
     vx_size pixelSizeInBitsNum;
     vx_color_space_e colorSpace;
     vx_channel_range_e channelRange;
     vx_size pixelSizeInBitsDenom;
 } AgoImageFormatDescription;
  
 typedef struct
 {
     vx_uint32 device_type; // shall be AGO_TARGET_AFFINITY_CPU or AGO_TARGET_AFFINITY_GPU
     vx_uint32 device_info; // reserved -- shall be initialized to ZERO and shall not be modified
     vx_uint32 group; // reserved -- shall be initialized to ZERO and shall not be modified
     vx_uint32 reserved; // reserved -- shall be initialized to ZERO and shall not be modified
 } AgoTargetAffinityInfo;
  
 typedef struct
 {
     vx_char macroName[256];
     vx_char *text;
 } AgoContextTextMacroInfo;
  
 typedef struct
 {
     vx_char *text;
     vx_uint32 num_ref;
     vx_reference *ref;
     vx_int32 dumpToConsole;
     void(VX_CALLBACK *data_registry_callback_f)(void *obj, vx_reference ref, const char *name, const char *app_params);
     void *data_registry_callback_obj;
 } AgoGraphImportInfo;
  
 typedef struct
 {
     vx_char fileName[256];
     vx_uint32 num_ref;
     vx_reference *ref;
     vx_char comment[64];
 } AgoGraphExportInfo;
  
 typedef struct
 {
     vx_uint64 kernel_enqueue;
     vx_uint64 kernel_wait;
     vx_uint64 buffer_read;
     vx_uint64 buffer_write;
 } AgoGraphPerfInternalInfo;
  
 typedef struct AgoNodeMergeRule_t
 {
     struct
     {
         vx_enum kernel_id;
         vx_uint32 arg_spec[AGO_MAX_PARAMS];
     } find[AGO_MERGE_RULE_MAX_FIND];
     struct
     {
         vx_enum kernel_id;
         vx_uint32 arg_spec[AGO_MAX_PARAMS];
     } replace[AGO_MERGE_RULE_MAX_REPLACE];
 } AgoNodeMergeRule;
  
 #ifdef __cplusplus
 typedef vx_status(VX_CALLBACK *amd_kernel_query_target_support_f)(vx_graph graph, vx_node node,
                                                                   vx_bool use_opencl_1_2,              // [input]  false: OpenCL driver is 2.0+; true: OpenCL driver is 1.2
                                                                   vx_uint32 &supported_target_affinity // [output] must be set to AGO_TARGET_AFFINITY_CPU or AGO_TARGET_AFFINITY_GPU or (AGO_TARGET_AFFINITY_CPU | AGO_TARGET_AFFINITY_GPU)
 );
  
 typedef vx_status(VX_CALLBACK *amd_kernel_opencl_codegen_callback_f)(
     vx_node node,                                // [input] node
     const vx_reference parameters[],             // [input] parameters
     vx_uint32 num,                               // [input] number of parameters
     bool opencl_load_function,                   // [input]  false: normal OpenCL kernel; true: reserved
     char opencl_kernel_function_name[64],        // [output] kernel_name for clCreateKernel()
     std::string &opencl_kernel_code,             // [output] string for clCreateProgramWithSource()
     std::string &opencl_build_options,           // [output] options for clBuildProgram()
     vx_uint32 &opencl_work_dim,                  // [output] work_dim for clEnqueueNDRangeKernel()
     vx_size opencl_global_work[],                // [output] global_work[] for clEnqueueNDRangeKernel()
     vx_size opencl_local_work[],                 // [output] local_work[] for clEnqueueNDRangeKernel()
     vx_uint32 &opencl_local_buffer_usage_mask,   // [output] reserved: must be ZERO
     vx_uint32 &opencl_local_buffer_size_in_bytes // [output] reserved: must be ZERO
 );
  
 typedef vx_status(VX_CALLBACK *amd_drama_add_node_f)(vx_node node, vx_enum kernel_id, vx_reference *paramList, vx_uint32 paramCount);
 typedef vx_status(VX_CALLBACK *amd_kernel_node_regen_callback_f)(vx_node node, amd_drama_add_node_f add_node_f, vx_bool &replace_original);
  
 typedef vx_status(VX_CALLBACK *amd_kernel_opencl_global_work_update_callback_f)(
     vx_node node,                     // [input] node
     const vx_reference parameters[],  // [input] parameters
     vx_uint32 num,                    // [input] number of parameters
     vx_uint32 opencl_work_dim,        // [input] work_dim for clEnqueueNDRangeKernel()
     vx_size opencl_global_work[],     // [output] global_work[] for clEnqueueNDRangeKernel()
     const vx_size opencl_local_work[] // [input] local_work[] for clEnqueueNDRangeKernel()
 );
  
 typedef vx_status(VX_CALLBACK *amd_kernel_gpu_buffer_update_callback_f)(
     vx_node node,                    // [input] node
     const vx_reference parameters[], // [input] parameters
     vx_uint32 num                    // [input] number of parameters
 );
  
 typedef struct
 {
     amd_kernel_gpu_buffer_update_callback_f gpu_buffer_update_callback_f;
     vx_uint32 gpu_buffer_update_param_index;
 } AgoKernelGpuBufferUpdateInfo;
  
 #if defined(AMD_FP16_SUPPORT)
 #include <half/half.hpp>
 typedef half_float::half vx_float16;
 #endif
 #endif
  
 #ifdef __cplusplus
 extern "C"
 {
 #endif
  
     VX_API_ENTRY vx_status VX_API_CALL vxGetReferenceName(vx_reference ref, vx_char name[], vx_size size);
  
     VX_API_ENTRY vx_status VX_API_CALL vxSetModuleInternalData(vx_context context, const vx_char *module, void *ptr, vx_size size);
  
     VX_API_ENTRY vx_status VX_API_CALL vxGetModuleInternalData(vx_context context, const vx_char *module, void **ptr, vx_size *size);
  
     VX_API_ENTRY vx_status VX_API_CALL vxSetModuleHandle(vx_node node, const vx_char *module, void *ptr);
  
     VX_API_ENTRY vx_status VX_API_CALL vxGetModuleHandle(vx_node node, const vx_char *module, void **ptr);
  
     VX_API_ENTRY vx_status VX_API_CALL vxSetContextImageFormatDescription(vx_context context, vx_df_image format, const AgoImageFormatDescription *desc);
  
     VX_API_ENTRY vx_status VX_API_CALL vxGetContextImageFormatDescription(vx_context context, vx_df_image format, AgoImageFormatDescription *desc);
  
     VX_API_ENTRY vx_status VX_API_CALL vxAliasTensor(vx_tensor tensorMaster, vx_size offset, vx_tensor tensor);
  
     VX_API_ENTRY vx_bool VX_API_CALL vxIsTensorAliased(vx_tensor tensorMaster, vx_size offset, vx_tensor tensor);
  
 #ifdef __cplusplus
 }
 #endif
  
 #endif