1 /* 2 * Copyright © Microsoft Corporation 3 * 4 * Permission is hereby granted, free of charge, to any person obtaining a 5 * copy of this software and associated documentation files (the "Software"), 6 * to deal in the Software without restriction, including without limitation 7 * the rights to use, copy, modify, merge, publish, distribute, sublicense, 8 * and/or sell copies of the Software, and to permit persons to whom the 9 * Software is furnished to do so, subject to the following conditions: 10 * 11 * The above copyright notice and this permission notice (including the next 12 * paragraph) shall be included in all copies or substantial portions of the 13 * Software. 14 * 15 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR 16 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, 17 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL 18 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER 19 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING 20 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS 21 * IN THE SOFTWARE. 22 */ 23 24 #ifndef CLC_COMPILER_H 25 #define CLC_COMPILER_H 26 27 #include "clc/clc.h" 28 29 #ifdef __cplusplus 30 extern "C" { 31 #endif 32 33 #define CLC_MAX_CONSTS 32 34 #define CLC_MAX_BINDINGS_PER_ARG 3 35 #define CLC_MAX_SAMPLERS 16 36 37 struct clc_printf_info { 38 unsigned num_args; 39 unsigned *arg_sizes; 40 char *str; 41 }; 42 43 struct clc_dxil_metadata { 44 struct { 45 unsigned offset; 46 unsigned size; 47 union { 48 struct { 49 unsigned buf_ids[CLC_MAX_BINDINGS_PER_ARG]; 50 unsigned num_buf_ids; 51 } image; 52 struct { 53 unsigned sampler_id; 54 } sampler; 55 struct { 56 unsigned buf_id; 57 } globconstptr; 58 struct { 59 unsigned sharedmem_offset; 60 } localptr; 61 }; 62 } *args; 63 unsigned kernel_inputs_cbv_id; 64 unsigned kernel_inputs_buf_size; 65 unsigned work_properties_cbv_id; 66 size_t num_uavs; 67 size_t num_srvs; 68 size_t num_samplers; 69 70 struct { 71 void *data; 72 size_t size; 73 unsigned uav_id; 74 } consts[CLC_MAX_CONSTS]; 75 size_t num_consts; 76 77 struct { 78 unsigned sampler_id; 79 unsigned addressing_mode; 80 unsigned normalized_coords; 81 unsigned filter_mode; 82 } const_samplers[CLC_MAX_SAMPLERS]; 83 size_t num_const_samplers; 84 size_t local_mem_size; 85 size_t priv_mem_size; 86 87 uint16_t local_size[3]; 88 uint16_t local_size_hint[3]; 89 90 struct { 91 unsigned info_count; 92 struct clc_printf_info *infos; 93 int uav_id; 94 } printf; 95 }; 96 97 struct clc_dxil_object { 98 const struct clc_kernel_info *kernel; 99 struct clc_dxil_metadata metadata; 100 struct { 101 void *data; 102 size_t size; 103 } binary; 104 }; 105 106 struct clc_runtime_arg_info { 107 union { 108 struct { 109 unsigned size; 110 } localptr; 111 struct { 112 unsigned normalized_coords; 113 unsigned addressing_mode; /* See SPIR-V spec for value meanings */ 114 unsigned linear_filtering; 115 } sampler; 116 }; 117 }; 118 119 struct clc_runtime_kernel_conf { 120 uint16_t local_size[3]; 121 struct clc_runtime_arg_info *args; 122 unsigned lower_bit_size; 123 unsigned support_global_work_id_offsets; 124 unsigned support_workgroup_id_offsets; 125 }; 126 127 struct clc_libclc_dxil_options { 128 unsigned optimize; 129 }; 130 131 struct clc_libclc * 132 clc_libclc_new_dxil(const struct clc_logger *logger, 133 const struct clc_libclc_dxil_options *dxil_options); 134 135 bool 136 clc_spirv_to_dxil(struct clc_libclc *lib, 137 const struct clc_binary *linked_spirv, 138 const struct clc_parsed_spirv *parsed_data, 139 const char *entrypoint, 140 const struct clc_runtime_kernel_conf *conf, 141 const struct clc_spirv_specialization_consts *consts, 142 const struct clc_logger *logger, 143 struct clc_dxil_object *out_dxil); 144 145 void clc_free_dxil_object(struct clc_dxil_object *dxil); 146 147 /* This struct describes the layout of data expected in the CB bound at global_work_offset_cbv_id */ 148 struct clc_work_properties_data { 149 /* Returned from get_global_offset(), and added into get_global_id() */ 150 unsigned global_offset_x; 151 unsigned global_offset_y; 152 unsigned global_offset_z; 153 /* Returned from get_work_dim() */ 154 unsigned work_dim; 155 /* The number of work groups being launched (i.e. the parameters to Dispatch). 156 * If the requested global size doesn't fit in a single Dispatch, these values should 157 * indicate the total number of groups that *should* have been launched. */ 158 unsigned group_count_total_x; 159 unsigned group_count_total_y; 160 unsigned group_count_total_z; 161 unsigned padding; 162 /* If the requested global size doesn't fit in a single Dispatch, subsequent dispatches 163 * should fill out these offsets to indicate how many groups have already been launched */ 164 unsigned group_id_offset_x; 165 unsigned group_id_offset_y; 166 unsigned group_id_offset_z; 167 }; 168 169 uint64_t clc_compiler_get_version(void); 170 171 #ifdef __cplusplus 172 } 173 #endif 174 175 #endif 176