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 #include "dxil_versions.h" 29 30 #ifdef __cplusplus 31 extern "C" { 32 #endif 33 34 #define CLC_MAX_CONSTS 32 35 #define CLC_MAX_BINDINGS_PER_ARG 3 36 #define CLC_MAX_SAMPLERS 16 37 38 struct clc_printf_info { 39 unsigned num_args; 40 unsigned *arg_sizes; 41 char *str; 42 }; 43 44 struct clc_dxil_metadata { 45 struct { 46 unsigned offset; 47 unsigned size; 48 union { 49 struct { 50 unsigned buf_ids[CLC_MAX_BINDINGS_PER_ARG]; 51 unsigned num_buf_ids; 52 } image; 53 struct { 54 unsigned sampler_id; 55 } sampler; 56 struct { 57 unsigned buf_id; 58 } globconstptr; 59 struct { 60 unsigned sharedmem_offset; 61 } localptr; 62 }; 63 } *args; 64 unsigned kernel_inputs_cbv_id; 65 unsigned kernel_inputs_buf_size; 66 unsigned work_properties_cbv_id; 67 size_t num_uavs; 68 size_t num_srvs; 69 size_t num_samplers; 70 71 struct { 72 void *data; 73 size_t size; 74 unsigned uav_id; 75 } consts[CLC_MAX_CONSTS]; 76 size_t num_consts; 77 78 struct { 79 unsigned sampler_id; 80 unsigned addressing_mode; 81 unsigned normalized_coords; 82 unsigned filter_mode; 83 } const_samplers[CLC_MAX_SAMPLERS]; 84 size_t num_const_samplers; 85 size_t local_mem_size; 86 size_t priv_mem_size; 87 88 uint16_t local_size[3]; 89 uint16_t local_size_hint[3]; 90 91 struct { 92 unsigned info_count; 93 struct clc_printf_info *infos; 94 int uav_id; 95 } printf; 96 }; 97 98 struct clc_dxil_object { 99 const struct clc_kernel_info *kernel; 100 struct clc_dxil_metadata metadata; 101 struct { 102 void *data; 103 size_t size; 104 } binary; 105 }; 106 107 struct clc_runtime_arg_info { 108 union { 109 struct { 110 unsigned size; 111 } localptr; 112 struct { 113 unsigned normalized_coords; 114 unsigned addressing_mode; /* See SPIR-V spec for value meanings */ 115 unsigned linear_filtering; 116 } sampler; 117 }; 118 }; 119 120 struct clc_runtime_kernel_conf { 121 uint16_t local_size[3]; 122 struct clc_runtime_arg_info *args; 123 unsigned lower_bit_size; 124 unsigned support_global_work_id_offsets; 125 unsigned support_workgroup_id_offsets; 126 127 enum dxil_shader_model max_shader_model; 128 enum dxil_validator_version validator_version; 129 }; 130 131 struct clc_libclc_dxil_options { 132 unsigned optimize; 133 }; 134 135 struct clc_libclc * 136 clc_libclc_new_dxil(const struct clc_logger *logger, 137 const struct clc_libclc_dxil_options *dxil_options); 138 139 bool 140 clc_spirv_to_dxil(struct clc_libclc *lib, 141 const struct clc_binary *linked_spirv, 142 const struct clc_parsed_spirv *parsed_data, 143 const char *entrypoint, 144 const struct clc_runtime_kernel_conf *conf, 145 const struct clc_spirv_specialization_consts *consts, 146 const struct clc_logger *logger, 147 struct clc_dxil_object *out_dxil); 148 149 void clc_free_dxil_object(struct clc_dxil_object *dxil); 150 151 /* This struct describes the layout of data expected in the CB bound at global_work_offset_cbv_id */ 152 struct clc_work_properties_data { 153 /* Returned from get_global_offset(), and added into get_global_id() */ 154 unsigned global_offset_x; 155 unsigned global_offset_y; 156 unsigned global_offset_z; 157 /* Returned from get_work_dim() */ 158 unsigned work_dim; 159 /* The number of work groups being launched (i.e. the parameters to Dispatch). 160 * If the requested global size doesn't fit in a single Dispatch, these values should 161 * indicate the total number of groups that *should* have been launched. */ 162 unsigned group_count_total_x; 163 unsigned group_count_total_y; 164 unsigned group_count_total_z; 165 unsigned padding; 166 /* If the requested global size doesn't fit in a single Dispatch, subsequent dispatches 167 * should fill out these offsets to indicate how many groups have already been launched */ 168 unsigned group_id_offset_x; 169 unsigned group_id_offset_y; 170 unsigned group_id_offset_z; 171 }; 172 173 uint64_t clc_compiler_get_version(void); 174 175 #ifdef __cplusplus 176 } 177 #endif 178 179 #endif 180