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();
170 
171 #ifdef __cplusplus
172 }
173 #endif
174 
175 #endif
176