mirror of https://gitlab.freedesktop.org/mesa/mesa
176 lines
5.0 KiB
C
176 lines
5.0 KiB
C
/*
|
||
* Copyright © Microsoft Corporation
|
||
*
|
||
* 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 (including the next
|
||
* paragraph) 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 CLC_COMPILER_H
|
||
#define CLC_COMPILER_H
|
||
|
||
#include "clc/clc.h"
|
||
|
||
#ifdef __cplusplus
|
||
extern "C" {
|
||
#endif
|
||
|
||
#define CLC_MAX_CONSTS 32
|
||
#define CLC_MAX_BINDINGS_PER_ARG 3
|
||
#define CLC_MAX_SAMPLERS 16
|
||
|
||
struct clc_printf_info {
|
||
unsigned num_args;
|
||
unsigned *arg_sizes;
|
||
char *str;
|
||
};
|
||
|
||
struct clc_dxil_metadata {
|
||
struct {
|
||
unsigned offset;
|
||
unsigned size;
|
||
union {
|
||
struct {
|
||
unsigned buf_ids[CLC_MAX_BINDINGS_PER_ARG];
|
||
unsigned num_buf_ids;
|
||
} image;
|
||
struct {
|
||
unsigned sampler_id;
|
||
} sampler;
|
||
struct {
|
||
unsigned buf_id;
|
||
} globconstptr;
|
||
struct {
|
||
unsigned sharedmem_offset;
|
||
} localptr;
|
||
};
|
||
} *args;
|
||
unsigned kernel_inputs_cbv_id;
|
||
unsigned kernel_inputs_buf_size;
|
||
unsigned work_properties_cbv_id;
|
||
size_t num_uavs;
|
||
size_t num_srvs;
|
||
size_t num_samplers;
|
||
|
||
struct {
|
||
void *data;
|
||
size_t size;
|
||
unsigned uav_id;
|
||
} consts[CLC_MAX_CONSTS];
|
||
size_t num_consts;
|
||
|
||
struct {
|
||
unsigned sampler_id;
|
||
unsigned addressing_mode;
|
||
unsigned normalized_coords;
|
||
unsigned filter_mode;
|
||
} const_samplers[CLC_MAX_SAMPLERS];
|
||
size_t num_const_samplers;
|
||
size_t local_mem_size;
|
||
size_t priv_mem_size;
|
||
|
||
uint16_t local_size[3];
|
||
uint16_t local_size_hint[3];
|
||
|
||
struct {
|
||
unsigned info_count;
|
||
struct clc_printf_info *infos;
|
||
int uav_id;
|
||
} printf;
|
||
};
|
||
|
||
struct clc_dxil_object {
|
||
const struct clc_kernel_info *kernel;
|
||
struct clc_dxil_metadata metadata;
|
||
struct {
|
||
void *data;
|
||
size_t size;
|
||
} binary;
|
||
};
|
||
|
||
struct clc_runtime_arg_info {
|
||
union {
|
||
struct {
|
||
unsigned size;
|
||
} localptr;
|
||
struct {
|
||
unsigned normalized_coords;
|
||
unsigned addressing_mode; /* See SPIR-V spec for value meanings */
|
||
unsigned linear_filtering;
|
||
} sampler;
|
||
};
|
||
};
|
||
|
||
struct clc_runtime_kernel_conf {
|
||
uint16_t local_size[3];
|
||
struct clc_runtime_arg_info *args;
|
||
unsigned lower_bit_size;
|
||
unsigned support_global_work_id_offsets;
|
||
unsigned support_workgroup_id_offsets;
|
||
};
|
||
|
||
struct clc_libclc_dxil_options {
|
||
unsigned optimize;
|
||
};
|
||
|
||
struct clc_libclc *
|
||
clc_libclc_new_dxil(const struct clc_logger *logger,
|
||
const struct clc_libclc_dxil_options *dxil_options);
|
||
|
||
bool
|
||
clc_spirv_to_dxil(struct clc_libclc *lib,
|
||
const struct clc_binary *linked_spirv,
|
||
const struct clc_parsed_spirv *parsed_data,
|
||
const char *entrypoint,
|
||
const struct clc_runtime_kernel_conf *conf,
|
||
const struct clc_spirv_specialization_consts *consts,
|
||
const struct clc_logger *logger,
|
||
struct clc_dxil_object *out_dxil);
|
||
|
||
void clc_free_dxil_object(struct clc_dxil_object *dxil);
|
||
|
||
/* This struct describes the layout of data expected in the CB bound at global_work_offset_cbv_id */
|
||
struct clc_work_properties_data {
|
||
/* Returned from get_global_offset(), and added into get_global_id() */
|
||
unsigned global_offset_x;
|
||
unsigned global_offset_y;
|
||
unsigned global_offset_z;
|
||
/* Returned from get_work_dim() */
|
||
unsigned work_dim;
|
||
/* The number of work groups being launched (i.e. the parameters to Dispatch).
|
||
* If the requested global size doesn't fit in a single Dispatch, these values should
|
||
* indicate the total number of groups that *should* have been launched. */
|
||
unsigned group_count_total_x;
|
||
unsigned group_count_total_y;
|
||
unsigned group_count_total_z;
|
||
unsigned padding;
|
||
/* If the requested global size doesn't fit in a single Dispatch, subsequent dispatches
|
||
* should fill out these offsets to indicate how many groups have already been launched */
|
||
unsigned group_id_offset_x;
|
||
unsigned group_id_offset_y;
|
||
unsigned group_id_offset_z;
|
||
};
|
||
|
||
uint64_t clc_compiler_get_version();
|
||
|
||
#ifdef __cplusplus
|
||
}
|
||
#endif
|
||
|
||
#endif
|