def prospero(src, dst):
print('--> compiling prospero-wave-psslc {}'.format(src))
bin = get_compiler_from_env(*_config[Languages.PROSPERO])
params = ['-DGNM', '-O4']
if '.frag' in src:
params += ['-profile', 'sce_ps_prospero']
if '.vert' in src:
params += ['-profile', 'sce_vs_vs_prospero']
if '.comp' in src:
params += ['-profile', 'sce_cs_prospero']
if '.tesc' in src:
params += ['-profile', 'sce_hs_prospero']
if '.tese' in src:
params += ['-profile', 'sce_ds_vs_prospero']
if '.geom' in src:
params += ['-profile', 'sce_gs_prospero']
params += ['-I' + fsl_basepath, '-o', dst, src]
status, output = get_status(bin, params)
fsl_assert(status == 0, src, message=output)
if output:
print(output)
return status
def d3d11(src, dst):
print('--> compiling DIRECT3D11 (fxc) {}'.format(src))
bin = get_compiler_from_env(*_config[Languages.DIRECT3D11])
params = ['/Zi']
if '.frag' in src:
params += ['/T', 'ps_5_0']
if '.vert' in src:
params += ['/T', 'vs_5_0']
if '.comp' in src:
params += ['/T', 'cs_5_0']
# handling network path on windows
if src.startswith("//"):
src = src.replace("//", "\\\\", 1)
if dst.startswith("//"):
dst = dst.replace("//", "\\\\", 1)
params += ['/I', fsl_basepath, '/Fo', dst, src]
# params = []
# params += ['/I', fsl_basepath, '/P', dst, src]
status, output = get_status(bin, params)
fsl_assert(status == 0, src, message=output)
return status
def vulkan(src, dst):
print('--> compiling VULKAN (glslangValidator) {}'.format(src))
bin = get_compiler_from_env(*_config[Languages.VULKAN])
''' prepend glsl directives '''
params = ['-V', src, '-o', dst, '-I' + fsl_basepath]
# params = ['-DFSL_GLSL', '-DSAMPLE_COUNT 2', '-V', src, '-o', dst, '-I'+fsl_basepath]
# params = ['-DFSL_GLSL', '-DSAMPLE_COUNT 2', src, '-I'+fsl_basepath, '-E']
if '.frag' in src:
params += ['-S', 'frag']
if '.vert' in src:
params += ['-S', 'vert']
if '.comp' in src:
params += ['-S', 'comp']
if '.geom' in src:
params += ['-S', 'geom']
if '.tesc' in src:
params += ['-S', 'tesc']
if '.tese' in src:
params += ['-S', 'tese']
params += ['--target-env', 'vulkan1.1']
status, output = get_status(bin, params)
fsl_assert(status == 0, src, message=output)
return status
def scarlett(src, dst):
print('--> compiling (dxc) SCARLETT {}'.format(src))
bin = get_compiler_from_env(*_config[Languages.SCARLETT])
params = ['/Zi', '-Qembed_debug']
if '.frag' in src:
params += ['/T', 'ps_6_4']
if '.vert' in src:
params += ['/T', 'vs_6_4']
if '.comp' in src:
params += ['/T', 'cs_6_0']
if '.tesc' in src:
params += ['/T', 'hs_6_0']
if '.tese' in src:
params += ['/T', 'ds_6_0']
params += [
'/I', fsl_basepath, '/D__XBOX_DISABLE_PRECOMPILE', '/Fo', dst, src
]
status, output = get_status(bin, params)
fsl_assert(status == 0, src, message=output)
if output:
print(output)
return status
return status
def orbis(src, dst):
print('--> compiling orbis-wave-psslc {}'.format(src))
bin = get_compiler_from_env(*_config[Languages.ORBIS])
params = ['-DGNM', '-O4']
if '.frag' in src:
params += ['-profile', 'sce_ps_orbis']
if '.vert' in src:
params += ['-profile', 'sce_vs_vs_orbis']
if '.comp' in src:
params += ['-profile', 'sce_cs_orbis']
if '.tesc' in src:
params += ['-profile', 'sce_hs_off_chip_orbis']
if '.tese' in src:
params += ['-profile', 'sce_ds_vs_off_chip_orbis']
if '.geom' in src:
params += ['-profile', 'sce_gs_orbis']
params += ['-I' + fsl_basepath, '-o', dst + '.sb', src]
# params += ['-I'+fsl_basepath, '-o', dst + '.sb', src, '-E']
# params += ['-I'+fsl_basepath, src, '-E']
status, output = get_status(bin, params)
fsl_assert(status == 0, src, message=output)
if output:
print(output)
return status
def get_mtl_patch_type(tessellation_layout):
mtl_patch_types = {
'quad': 'quad',
'triangle': 'triangle',
}
domain = tessellation_layout[0].strip('"')
fsl_assert(domain in mtl_patch_types, message='Cannot map domain to mtl patch type: ' + domain)
return mtl_patch_types[domain]
def metal(src, dst):
print('--> compiling METAL (metal.exe) {}'.format(src))
bin = get_compiler_from_env(*_config[Languages.METAL])
# params = ['-D','FSL_METAL', '-dD', '-C', '-E', '-o', dst, tmp]
# params = ['-D','FSL_METAL', '-dD', '-S', '-o', dst, tmp]
# params = ['-dD', '-I', fsl_basepath, '-o', dst, tmp]
params = ['-I', fsl_basepath]
params += ['-dD', '-o', dst, src]
status, output = get_status(bin, params)
fsl_assert(status == 0, src, message=output)
return status
def gles(src, dst):
print('--> compiling GLES using (glslangValidator) {}'.format(src))
bin = get_compiler_from_env(*_config[Languages.GLES])
''' prepend glsl directives '''
params = [src, '-I' + fsl_basepath]
if '.frag' in src:
params += ['-S', 'frag']
if '.vert' in src:
params += ['-S', 'vert']
if '.comp' in src:
params += ['-S', 'comp']
if '.geom' in src:
params += ['-S', 'geom']
if '.tesc' in src:
params += ['-S', 'tesc']
if '.tese' in src:
params += ['-S', 'tese']
status, output = get_status(bin, params)
fsl_assert(status == 0, src, message=output)
return status
def d3d12(src, dst):
print('--> compiling DIRECT3D12 (dxc) {}'.format(src))
bin = get_compiler_from_env(*_config[Languages.DIRECT3D12])
params = []
if '.frag' in src:
params += ['/T', 'ps_6_4']
if '.vert' in src:
params += ['/T', 'vs_6_4']
if '.comp' in src:
params += ['/T', 'cs_6_0']
if '.geom' in src:
params += ['/T', 'gs_6_0']
if '.tesc' in src:
params += ['/T', 'hs_6_0']
if '.tese' in src:
params += ['/T', 'ds_6_0']
# params += ['/I', fsl_basepath, '/DFSL_D3D12', '/Fo', dst, src]
params += ['/I', fsl_basepath, '/Fo', dst, src]
status, output = get_status(bin, params)
fsl_assert(status == 0, src, message=output)
return status
def vulkan(fsl, dst):
shader = getShader(fsl, dst)
shader_src = getHeader(fsl)
shader_src += [
'#version 450 core\n',
'#extension GL_GOOGLE_include_directive : require\nprecision highp float;\nprecision highp int;\n\n'
]
shader_src += ['#define STAGE_', shader.stage.name, '\n']
if shader.enable_waveops:
shader_src += ['#define ENABLE_WAVEOPS()\n']
in_location = 0
# shader_src += ['#include "includes/vulkan.h"\n\n']
# incPath = os.path.join(os.path.dirname(os.path.dirname(__file__)), 'includes', 'vulkan.h')
# dstInc = os.path.join(os.path.dirname(dst), "includes/vulkan.h")
# if True or not os.path.exists(dstInc):
# os.makedirs(os.path.dirname(dstInc), exist_ok=True)
# copyfile(incPath, dstInc)
# directly embed vk header in shader
header_path = os.path.join(os.path.dirname(os.path.dirname(__file__)),
'includes', 'vulkan.h')
header_lines = open(header_path).readlines()
shader_src += header_lines + ['\n']
# pcf output defines (inputs are identical to main)
pcf_return_assignments = []
# if shader.stage == Stages.TESC and shader.pcf:
# t = getMacroName(shader.pcf_returnType)
# if t in shader.structs:
# for t, n, s in shader.structs[t]:
# pcf_return_assignments += ['ASSIGN_PCF_OUT_' + getArrayBaseName(n)]
# retrieve output patch size
patch_size = 0
for line in shader.lines:
if 'OUTPUT_CONTROL_POINTS' in line:
patch_size = int(getMacro(line))
out_location = 0
arrBuffs = [
'Get(' + getArrayBaseName(res[1]) + ')' for res in shader.resources
if 'Buffer' in res[0] and (isArray(res[1]))
]
returnType = None if not shader.returnType else (getMacroName(
shader.returnType), getMacro(shader.returnType))
push_constant = None
skip_line = False
parsing_struct = None
parsing_cbuffer = None
parsing_pushconstant = None
nonuniformresourceindex = None
parsed_entry = False
# tesselation
pcf_returnType = None
pcf_arguments = []
substitutions = {
'min16float(': 'float(',
'min16float2(': 'float2(',
'min16float3(': 'float3(',
'min16float4(': 'float4(',
}
if shader.returnType and shader.returnType in shader.structs:
for _, _, sem in shader.structs[shader.returnType]:
sem = sem.upper()
if sem == 'SV_RENDERTARGETARRAYINDEX':
shader_src += [
'#extension GL_ARB_shader_viewport_layer_array : enable\n\n'
]
break
struct_declarations = []
input_assignments = []
return_assignments = []
shader_src += ['#line 1 \"' + fsl.replace(os.sep, '/') + '\"\n']
line_index = 0
for line in shader.lines:
line_index += 1
shader_src_len = len(shader_src)
if line.startswith('#line'):
line_index = int(line.split()[1]) - 1
def get_uid(name):
return '_' + name + '_' + str(len(shader_src))
# dont process commented lines
if line.strip().startswith('//'):
shader_src += [line]
continue
# TODO: handle this differently
if '#ifdef NO_FSL_DEFINITIONS' in line:
skip_line = True
if skip_line and '#endif' in line:
skip_line = False
continue
if skip_line:
continue
for k, v in substitutions.items():
l0 = line.find(k)
while l0 > -1:
line = line.replace(k, v)
l0 = line.find(k)
if line.strip().startswith('STRUCT('):
parsing_struct = getMacro(line)
if parsing_struct and '};' in line:
if not line.endswith('\n'): line += '\n'
shader_src += [line]
for macro, struct_declaration in struct_declarations:
shader_src += ['#ifdef ' + macro + '\n']
shader_src += [''.join(struct_declaration) + '\n']
shader_src += ['#endif\n']
shader_src += ['#line {}\n'.format(line_index + 1)]
struct_declarations = []
parsing_struct = None
continue
# handle struct data declarations
if parsing_struct and line.strip().startswith('DATA('):
# handle entry return type
if shader.returnType and parsing_struct in shader.returnType:
var = 'out_' + shader.returnType
elem_dtype, elem_name, sem = getMacro(line)
sem = sem.upper()
flat_modifier = 'FLAT(' in line
if flat_modifier:
elem_dtype = getMacro(elem_dtype)
line = get_whitespace(line) + getMacro(
elem_dtype) + ' ' + elem_name + ';\n'
basename = getArrayBaseName(elem_name)
macro = get_uid(basename)
shader_src += ['#define ', macro, '\n']
output_datapath = var + '_' + elem_name
reference = None
if sem == 'SV_POSITION':
output_datapath = 'gl_Position'
if shader.stage == Stages.TESC:
output_datapath = 'gl_out[gl_InvocationID].' + output_datapath
elif sem == 'SV_POINTSIZE':
output_datapath = 'gl_PointSize'
elif sem == 'SV_DEPTH':
output_datapath = 'gl_FragDepth'
elif sem == 'SV_RENDERTARGETARRAYINDEX':
output_datapath = 'gl_Layer'
else:
output_prefix, out_postfix = '', ''
if shader.stage == Stages.TESC:
output_prefix = 'patch '
out_postfix = '[]'
if shader.input_patch_arg:
out_postfix = '[' + shader.input_patch_arg[1] + ']'
if flat_modifier:
output_prefix = 'flat ' + output_prefix
reference = [
'layout(location = ',
str(out_location), ') ', output_prefix, 'out(',
elem_dtype, ') ', output_datapath, out_postfix, ';'
]
if shader.stage == Stages.TESC and sem != 'SV_POSITION':
output_datapath += '[gl_InvocationID]'
# unroll attribute arrays
if isArray(elem_name):
# assignment = ['for(int i=0; i<', str(getArrayLenFlat(elem_name)), ';i++) ']
# assignment += [var, '_', basename, '[i] = ', var, '.', basename, '[i]']
assignment = [
var, '_', basename, ' = ', var, '.', basename, ''
]
out_location += getArrayLen(shader.defines, elem_name)
else:
if sem != 'SV_POSITION' and sem != 'SV_POINTSIZE' and sem != 'SV_DEPTH' and sem != 'SV_RENDERTARGETARRAYINDEX':
out_location += 1
if output_datapath == 'gl_Layer':
assignment = [
output_datapath, ' = int(', var, '.', elem_name,
')'
]
else:
assignment = [
output_datapath, ' = ', var, '.', elem_name
]
if reference:
struct_declarations += [(macro, reference)]
return_assignments += [(macro, assignment)]
# tesselation pcf output
elif shader.pcf_returnType and parsing_struct in shader.pcf_returnType:
# var = getMacro(shader.pcf_returnType)
var = 'out_' + shader.pcf_returnType
_, elem_name, sem = getMacro(line)
sem = sem.upper()
basename = getArrayBaseName(elem_name)
macro = get_uid(basename)
shader_src += ['#define ', macro, '\n']
tess_var = ''
if sem == 'SV_TESSFACTOR':
tess_var = 'gl_TessLevelOuter'
if sem == 'SV_INSIDETESSFACTOR':
tess_var = 'gl_TessLevelInner'
pcf_return_assignments += [
(macro, [tess_var, ' = ', var, '.', basename, ';'])
]
# elif shader.entry_arg and parsing_struct in shader.entry_arg:
elif is_input_struct(parsing_struct, shader):
var = get_input_struct_var(parsing_struct, shader)
elem_dtype, elem_name, sem = getMacro(line)
sem = sem.upper()
flat_modifier = 'FLAT(' in line
if flat_modifier:
elem_dtype = getMacro(elem_dtype)
line = get_whitespace(line) + getMacro(
elem_dtype) + ' ' + elem_name + ';\n'
is_array = isArray(elem_name)
basename = getArrayBaseName(elem_name)
macro = get_uid(basename)
shader_src += ['#define ', macro, '\n']
input_datapath = var + '_' + elem_name
if sem == 'SV_POINTSIZE' or sem == 'SV_RENDERTARGETARRAYINDEX':
continue
# for vertex shaders, use the semantic as attribute name (for name-based reflection)
input_value = sem if shader.stage == Stages.VERT else var + '_' + elem_name
# tessellation
in_postfix, in_prefix = '', ''
if shader.stage == Stages.TESC: in_postfix = '[]'
if shader.stage == Stages.TESE:
in_prefix = ' patch'
in_postfix = '[]'
if shader.output_patch_arg:
in_postfix = '[' + shader.output_patch_arg[1] + ']'
if flat_modifier:
in_prefix = 'flat ' + in_prefix
reference = [
'layout(location = ',
str(in_location), ')', in_prefix, ' in(', elem_dtype, ') ',
input_value, in_postfix, ';'
]
# input semantics
if sem == 'SV_POSITION' and shader.stage == Stages.FRAG:
input_value = elem_dtype + '(float4(gl_FragCoord.xyz, 1.0f / gl_FragCoord.w))'
reference = []
var_postfix = ''
if shader.stage == Stages.TESC:
if sem == 'SV_POSITION':
input_value = 'gl_in[gl_InvocationID].gl_Position'
reference = []
else:
input_value = input_value + '[gl_InvocationID]'
var_postfix = '[gl_InvocationID]'
if shader.stage == Stages.TESE:
if sem == 'SV_POSITION':
input_value = 'gl_in[0].gl_Position'
reference = []
elif shader.output_patch_arg and shader.output_patch_arg[
0] in parsing_struct:
input_value = input_value + '[0]'
var_postfix = '[0]'
if sem == 'SV_TESSFACTOR':
input_value = 'gl_TessLevelOuter'
var_postfix = ''
reference = []
if sem == 'SV_INSIDETESSFACTOR':
input_value = 'gl_TessLevelInner'
var_postfix = ''
reference = []
assignment = []
if sem == 'SV_VERTEXID': input_value = 'gl_VertexIndex'
# unroll attribute arrays
# if is_array:
# assignment = ['for(int i=0; i<', str(getArrayLenFlat(elem_name)), ';i++) ']
# assignment += [var, '.', basename, '[i] = ', var, '_', basename, '[i]']
# else:
assignment = [
var, var_postfix, '.', basename, ' = ', input_value
]
if shader.stage == Stages.VERT and sem != 'SV_VERTEXID':
assignment += [';\n\t', sem]
if sem != 'SV_POSITION':
in_location += 1
if reference:
struct_declarations += [(macro, reference)]
input_assignments += [(macro, assignment)]
if (parsing_cbuffer
or parsing_pushconstant) and line.strip().startswith('DATA('):
dt, name, sem = getMacro(line)
element_basename = getArrayBaseName(name)
element_path = None
if parsing_cbuffer:
element_path = element_basename
if parsing_pushconstant:
element_path = parsing_pushconstant[0] + '.' + element_basename
shader_src += [
'#define _Get', element_basename, ' ', element_path, '\n'
]
if 'CBUFFER' in line:
fsl_assert(parsing_cbuffer == None,
message='Inconsistent cbuffer declaration: \"' + line +
'\"')
parsing_cbuffer = tuple(getMacro(line))
if '};' in line and parsing_cbuffer:
parsing_cbuffer = None
if 'PUSH_CONSTANT' in line:
parsing_pushconstant = tuple(getMacro(line))
if '};' in line and parsing_pushconstant:
parsing_pushconstant = None
if is_groupshared_decl(line):
dtype, dname = getMacro(line)
basename = getArrayBaseName(dname)
shader_src += ['#define _Get', basename, ' ', basename, '\n']
line = 'shared ' + dtype + ' ' + dname + ';\n'
if 'EARLY_FRAGMENT_TESTS' in line:
line = 'layout(early_fragment_tests) in;\n'
if 'EnablePSInterlock' in line:
line = '#ifdef GL_ARB_fragment_shader_interlock\n' \
'layout(pixel_interlock_ordered) in;\n' \
'#endif\n'
# handle push constants
if line.strip().startswith('PUSH_CONSTANT'):
# push_constant = getMacro(line)[0]
push_constant = tuple(getMacro(line))
if push_constant and '};' in line:
shader_src += ['} ', push_constant[0], ';\n']
for dt, dn, _ in shader.pushConstant[push_constant]:
dn = getArrayBaseName(dn)
push_constant = None
continue
resource_decl = None
if line.strip().startswith('RES('):
resource_decl = getMacro(line)
fsl_assert(len(resource_decl) == 5,
fsl,
message='invalid Res declaration: \'' + line + '\'')
basename = getArrayBaseName(resource_decl[1])
shader_src += ['#define _Get' + basename + ' ' + basename + '\n']
# handle buffer resource declarations
if resource_decl and is_buffer(resource_decl):
shader_src += declare_buffer(resource_decl)
continue
# handle references to arrays of structured buffers
for buffer_array in arrBuffs:
line = insert_buffer_array_indirections(line, buffer_array)
# specify format qualified image resource
if resource_decl and is_rw_texture(resource_decl):
shader_src += declare_rw_texture(resource_decl)
continue
if 'BeginNonUniformResourceIndex(' in line:
index, max_index = getMacro(line), None
assert index != [], 'No index provided for {}'.format(line)
if type(index) == list:
max_index = index[1]
index = index[0]
nonuniformresourceindex = index
# if type(max_index) is str:
if max_index and not max_index.isnumeric():
max_index = ''.join(c for c in shader.defines[max_index]
if c.isdigit())
shader_src += BeginNonUniformResourceIndex(nonuniformresourceindex,
max_index)
continue
if 'EndNonUniformResourceIndex()' in line:
assert nonuniformresourceindex, 'EndNonUniformResourceIndex: BeginNonUniformResourceIndex not called/found'
shader_src += EndNonUniformResourceIndex(nonuniformresourceindex)
nonuniformresourceindex = None
continue
if nonuniformresourceindex:
shader_src += [line[:-1], ' \\\n']
continue
if '_MAIN(' in line:
if shader.returnType and shader.returnType not in shader.structs:
shader_src += [
'layout(location = 0) out(', shader.returnType, ') out_',
shader.returnType, ';\n'
]
if shader.input_patch_arg:
patch_size = shader.input_patch_arg[1]
shader_src += ['layout(vertices = ', patch_size, ') out;\n']
shader_src += ['void main()\n']
shader_src += ['#line {}\n'.format(line_index), '//' + line]
parsed_entry = True
continue
if parsed_entry and re.search('(^|\s+)RETURN', line):
ws = get_whitespace(line)
output_statement = [ws + '{\n']
if shader.returnType:
output_value = getMacro(line)
if shader.returnType not in shader.structs:
output_statement += [
ws + '\tout_' + shader.returnType + ' = ' +
output_value + ';\n'
]
else:
output_statement += [
ws + '\t' + shader.returnType + ' out_' +
shader.returnType + ' = ' + output_value + ';\n'
]
for macro, assignment in return_assignments:
output_statement += ['#ifdef ' + macro + '\n']
output_statement += [ws + '\t' + ''.join(assignment) + ';\n']
output_statement += ['#endif\n']
if shader.stage == Stages.TESC:
output_statement += ['\t\t' + shader.pcf + '();\n']
output_statement += [ws + '\treturn;\n' + ws + '}\n']
shader_src += output_statement
shader_src += ['#line {}\n'.format(line_index), '//' + line]
continue
if 'INIT_MAIN' in line:
# assemble input
for dtype, var in shader.struct_args:
if shader.input_patch_arg and dtype in shader.input_patch_arg[
0]:
dtype, dim, var = shader.input_patch_arg
shader_src += [
'\t' + dtype + ' ' + var + '[' + dim + '];\n'
]
continue
if shader.output_patch_arg and dtype in shader.output_patch_arg[
0]:
dtype, dim, var = shader.output_patch_arg
shader_src += [
'\t' + dtype + ' ' + var + '[' + dim + '];\n'
]
continue
shader_src += ['\t' + dtype + ' ' + var + ';\n']
for macro, assignment in input_assignments:
shader_src += ['#ifdef ' + macro + '\n']
shader_src += ['\t' + ''.join(assignment) + ';\n']
shader_src += ['#endif\n']
''' additional inputs '''
for dtype, dvar in shader.flat_args:
innertype = getMacro(dtype)
semtype = getMacroName(dtype)
shader_src += [
'\tconst ' + innertype + ' ' + dvar + ' = ' + innertype +
'(' + semtype.upper() + ');\n'
]
''' generate a statement for each vertex attribute
this should not be necessary, but it influences
the attribute order for the SpirVTools shader reflection
'''
# if shader.stage == Stages.VERT:# and shader.entry_arg:
# for dtype, var in shader.struct_args:
# for _, n, s in shader.structs[dtype]:
# if s.upper() == 'SV_VERTEXID': continue
# shader_src += ['\t', s ,';\n']
shader_src += ['#line {}\n'.format(line_index), '//' + line]
continue
# tesselation
if shader.pcf and shader.pcf in line and not pcf_returnType:
loc = line.find(shader.pcf)
pcf_returnType = line[:loc].strip()
pcf_arguments = getMacro(line[loc:])
_pcf_arguments = [
arg for arg in pcf_arguments if 'INPUT_PATCH' in arg
]
ws = line[:len(line) - len(line.lstrip())]
shader_src += [ws + 'void ' + shader.pcf + '()\n']
continue
if pcf_returnType and 'PCF_INIT' in line:
ws = line[:len(line) - len(line.lstrip())]
for dtype, dvar in shader.pcf_arguments:
sem = getMacroName(dtype)
innertype = getMacro(dtype)
print(innertype, sem, dvar)
if 'INPUT_PATCH' in sem:
shader_src += [ws + dtype + ' ' + dvar + ';\n']
else:
shader_src += [
ws + innertype + ' ' + dvar + ' = ' + sem.upper() +
';\n'
]
for macro, assignment in input_assignments:
shader_src += ['#ifdef ' + macro + '\n']
shader_src += ['\t' + ''.join(assignment) + ';\n']
shader_src += ['#endif\n']
shader_src += ['#line {}\n'.format(line_index), '//' + line]
continue
if pcf_returnType and 'PCF_RETURN' in line:
ws = get_whitespace(line)
output_statement = [ws + '{\n']
output_value = getMacro(line)
output_statement += [
ws + '\t' + pcf_returnType + ' out_' + pcf_returnType + ' = ' +
output_value + ';\n'
]
for macro, assignment in pcf_return_assignments:
output_statement += ['#ifdef ' + macro + '\n']
output_statement += [ws + '\t' + ''.join(assignment) + '\n']
output_statement += ['#endif\n']
output_statement += [ws + '\treturn;\n', ws + '}\n']
shader_src += output_statement
shader_src += ['#line {}\n'.format(line_index), '//' + line]
continue
if shader_src_len != len(shader_src):
shader_src += ['#line {}\n'.format(line_index)]
shader_src += [line]
open(dst, 'w').writelines(shader_src)
return 0
def metal(fsl, dst):
shader = getShader(fsl, dst)
msl_target = targetToMslEntry[shader.stage]
shader_src = getHeader(fsl)
shader_src += ['#define METAL\n']
if shader.enable_waveops:
shader_src += ['#define ENABLE_WAVEOPS()\n']
# shader_src += ['#include "includes/metal.h"\n']
# incPath = os.path.join(os.path.dirname(os.path.dirname(__file__)), 'includes', 'metal.h')
# dstInc = os.path.join(os.path.dirname(dst), "includes/metal.h")
# if True or not os.path.exists(dstInc):
# os.makedirs(os.path.dirname(dstInc), exist_ok=True)
# copyfile(incPath, dstInc)
# directly embed metal header in shader
header_path = os.path.join(os.path.dirname(os.path.dirname(__file__)), 'includes', 'metal.h')
header_lines = open(header_path).readlines()
shader_src += header_lines + ['\n']
def getMetalResourceID(dtype: str, reg: str) -> str:
S_OFFSET, U_OFFSET, B_OFFSET = 1024, 1024*2, 1024*3
if 'FSL_REG' in reg:
reg = getMacro(reg)[-1]
register = int(reg[1:])
if 't' == reg[0]: # SRV
return register
if 's' == reg[0]: # samplers
return S_OFFSET + register
if 'u' == reg[0]: # UAV
return U_OFFSET + register
if 'b' == reg[0]: # CBV
return B_OFFSET + register
assert False, '{}'.format(reg)
# for metal only consider resource declared using the following frequencies
metal_ab_frequencies = [
'UPDATE_FREQ_NONE',
'UPDATE_FREQ_PER_FRAME',
'UPDATE_FREQ_PER_BATCH',
'UPDATE_FREQ_PER_DRAW',
]
# collect array resources
ab_elements = {}
for resource_decl in shader.resources:
name, freq = resource_decl[1:3]
if isArray(name):
ab_elements[freq] = []
# tessellation layout, used for mtl patch type declaration
tessellation_layout = None
# consume any structured inputs, since an intermediate struct will be used and manually fed to DS
if shader.stage == Stages.TESE:
shader.struct_args = []
global_references = {}
global_reference_paths = {}
global_reference_args = {}
global_fn_table = get_fn_table(shader.lines)
# transform VS entry into callable fn
if shader.stage == Stages.TESC:
vertex_shader = getShader(shader.vs_reference)
vs_main, vs_parsing_main = [], -2
struct = None
elem_index = 0
# add vertex input as a shader resource to tesc
for struct in vertex_shader.struct_args:
res_decl = 'RES(Buffer(VSIn), vertexInput, UPDATE_FREQ_NONE, b0, binding = 0)'
shader.lines.insert(0, res_decl + ';\n')
shader.lines.insert(0, 'struct VSIn;\n')
# add mtl tessellation buffer outputs
r0_decl = 'RES(RWBuffer(HullOut), hullOutputBuffer, UPDATE_FREQ_NONE, b1, binding = 0)'
r0 = getMacroName(shader.returnType)
shader.lines.insert(0, r0_decl + ';\n')
shader.lines.insert(0, 'struct ' +r0+ ';\n')
r1_decl = 'RES(RWBuffer(PatchTess), tessellationFactorBuffer, UPDATE_FREQ_NONE, b2, binding = 0)'
r1 = getMacroName(shader.pcf_returnType)
shader.lines.insert(0, r1_decl + ';\n')
shader.lines.insert(0, 'struct ' +r1+ ';\n')
# TODO: this one is necessary to determine the total # of vertices to fetch from the buffer,
# this should be handled generically
di_decl = 'RES(Buffer(DrawIndirectInfo), drawInfo, UPDATE_FREQ_NONE, b3, binding = 0)'
shader.lines.insert(0, di_decl + ';\n')
shader.lines.insert(0, 'struct DrawIndirectInfo {uint vertexCount; uint _data[3];};\n')
# collect VSMain lines, transform inputs into const refs which will get manually fed from buffer
for line in vertex_shader.lines:
if line.strip().startswith('//'): continue
if line.strip().startswith('STRUCT('):
struct = getMacro(line)
if is_input_struct(struct, vertex_shader):
if line.strip().startswith('DATA('):
dtype, dname, _ = getMacro(line)
vs_main += [get_whitespace(line), dtype, ' ', dname, ' [[attribute(', str(elem_index), ')]];\n']
elem_index += 1
else:
vs_main += [line]
if struct and '};' in line:
struct = None
if 'VS_MAIN(' in line:
sig = line.strip().split(' ', 1)
line = getMacroName(sig[0]) + ' VSMain('
l0 = len(line)
prefix = ''
for dtype, dname in vertex_shader.struct_args:
line = prefix + line + 'constant ' + dtype + '& ' + dname
prefix = ', '
vs_parsing_main = -1
line = line+')\n'
vs_main += [line, '{\n']
global_fn_table['VSMain'] = (line, (999999, l0))
continue
if vs_parsing_main > -1:
l_get = line.find('Get(')
while l_get > 0:
resource = line[l_get:]
resource = resource[4:resource.find(')')]
fn = 'VSMain'
if fn not in global_references: global_references[fn] = set()
global_references[fn].add(resource)
l_get = line.find('Get(', l_get+1)
if 'INIT_MAIN' in line and vertex_shader.returnType:
continue
# mName = getMacroName(vertex_shader.returnType)
# mArg = getMacro(vertex_shader.returnType)
# line = get_whitespace(line) + '{} {};\n'.format(mName, mArg)
if re.search('(^|\s+)RETURN', line) and vertex_shader.returnType:
line = line.replace('RETURN', 'return ')
vs_main += [line]
if vs_parsing_main > -2 and '{' in line: vs_parsing_main += 1
if vs_parsing_main > -2 and '}' in line: vs_parsing_main -= 1
skip = False
struct = None
mainArgs = []
# statements which need to be inserted into INIT_MAIN
entry_declarations = []
parsing_cbuffer = None
struct_uid = None
parsing_pushconstant = None
# reserve the first 6 buffer locations for ABs(4) + rootcbv + rootconstant
buffer_location = 6
texture_location = 0
sampler_location = 0
attribute_index = 0
# reversed list of functions, used to determine where a resource is being accessed
reversed_fns = list(reversed(list(global_fn_table.items())))
parsing_main = False
global_scope_count = 0
for i, line in enumerate(shader.lines):
if line.strip().startswith('//'): continue
if '{' in line: global_scope_count += 1
if '}' in line: global_scope_count -= 1
l_get = line.find('Get(')
while l_get > 0 and not parsing_main:
resource = line[l_get:]
resource = resource[4:resource.find(')')]
for fn, (_, fn_i) in reversed_fns:
if fn_i[0] < i:
if fn not in global_references: global_references[fn] = set()
global_references[fn].add(resource)
break
l_get = line.find('Get(', l_get+1)
l_get = line.find('WaveGetLaneIndex()')
while l_get > 0 and not parsing_main:
resource = 'simd_lane_id'
for fn, (_, fn_i) in reversed_fns:
if fn_i[0] < i:
if fn not in global_references: global_references[fn] = set()
global_references[fn].add(resource)
break
l_get = line.find('Get(', l_get+1)
if not parsing_main:
global_references_tmp = list(global_references.items())
for fn, resource in global_references_tmp:
l_call = line.find(fn+'(')
if l_call > 0:
for fn_caller, (_, fn_i) in reversed_fns:
if fn_i[0] < i:
if fn_caller not in global_references: global_references[fn_caller] = set()
global_references[fn_caller].update(resource)
break
if '_MAIN(' in line:
parsing_main = True
if shader.enable_waveops:
global_reference_paths['simd_lane_id'] = 'simd_lane_id'
global_reference_args['simd_lane_id'] = 'const uint simd_lane_id'
def declare_argument_buffers(mainArgs):
ab_decl = []
# declare argument buffer structs
for freq, elements in ab_elements.items():
argBufType = 'AB_' + freq
ab_decl += ['\n\t// Generated Metal Resource Declaration: ', argBufType, '\n' ]
# skip empty update frequencies
if not elements: continue
# make AB declaration only active if any member is defined
resource_conditions = ' || '.join('defined('+macro+')' for macro, elem in elements)
ab_decl += ['#if ', resource_conditions , '\n']
ab_macro = get_uid(argBufType)
ab_decl += ['\t#define ', ab_macro, '\n']
space = 'constant'
mainArgs += [(ab_macro, [space, ' struct ', argBufType, '& ', argBufType, '[[buffer({})]]'.format(freq)])]
ab_decl += ['\tstruct ', argBufType, '\n\t{\n']
for macro, elem in elements:
ab_decl += ['\t\t#ifdef ', macro, '\n']
ab_decl += ['\t\t', *elem, ';\n']
ab_decl += ['\t\t#endif\n']
ab_decl += ['\t};\n']
ab_decl += ['#endif // End of Resource Declaration: ', argBufType, '\n']
return ab_decl
shader_src += ['#line 1 \"'+fsl.replace(os.sep, '/')+'\"\n']
line_index = 0
last_res_decl = 0
explicit_res_decl = None
for line in shader.lines:
line_index += 1
shader_src_len = len(shader_src)
if line.startswith('#line'):
line_index = int(line.split()[1]) - 1
def get_uid(name):
return name + '_' + str(len(shader_src))
# dont process commented lines
if line.strip().startswith('//'):
shader_src += ['\t', line]
continue
if 'DECLARE_RESOURCES' in line:
explicit_res_decl = len(shader_src) + 1
line = '//' + line
# TODO: improve this
if '#ifdef NO_FSL_DEFINITIONS' in line:
skip = True
if skip and '#endif' in line:
skip = False
continue
if skip:
continue
if line.strip().startswith('STRUCT('):
struct = getMacro(line)
if struct and '};' in line:
struct = None
if 'EARLY_FRAGMENT_TESTS' in line:
line = '[[early_fragment_tests]]\n'
if 'BeginNonUniformResourceIndex' in line:
nuri = getMacro(line)
if type(nuri) is str:
line = line.replace(nuri, nuri + ', None')
if 'TESS_LAYOUT(' in line:
tessellation_layout = getMacro(line)
# Shader I/O
if struct and line.strip().startswith('DATA('):
if shader.returnType and struct in shader.returnType:
var = getMacro(shader.returnType)
dtype, name, sem = getMacro(line)
sem = sem.upper()
base_name = getArrayBaseName(name)
macro = get_uid(base_name)
shader_src += ['#define ', macro, '\n']
output_semantic = ''
if 'SV_POSITION' in sem:
output_semantic = '[[position]]'
if 'SV_POINTSIZE' in sem:
output_semantic = '[[point_size]]'
if 'SV_DEPTH' in sem:
output_semantic = '[[depth(any)]]'
if 'SV_RENDERTARGETARRAYINDEX' in sem:
output_semantic = '[[render_target_array_index]]'
color_location = sem.find('SV_TARGET')
if color_location > -1:
color_location = sem[color_location+len('SV_TARGET'):]
if not color_location: color_location = '0'
output_semantic = '[[color('+color_location+')]]'
shader_src += ['#line {}\n'.format(line_index)]
shader_src += [get_whitespace(line), dtype, ' ', name, ' ', output_semantic, ';\n']
continue
elif is_input_struct(struct, shader):
var = get_input_struct_var(struct, shader)
dtype, name, sem = getMacro(line)
sem = sem.upper()
macro = get_uid(getArrayBaseName(name))
shader_src += ['#define ', macro, '\n']
# for vertex shaders, use the semantic as attribute name (for name-based reflection)
n2 = sem if shader.stage == Stages.VERT else getArrayBaseName(name)
if isArray(name):
base_name = getArrayBaseName(name)
array_length = int(getArrayLen(shader.defines, name))
assignment = []
for i in range(array_length): assignment += ['\t_',var, '.', base_name, '[', str(i), '] = ', var, '.', base_name, '_', str(i), '; \\\n']
# TODO: handle this case
attribute = ''
if shader.stage == Stages.VERT:
attribute = '[[attribute('+str(attribute_index)+')]]'
attribute_index += 1
elif shader.stage == Stages.FRAG:
attribute = ''
if 'SV_POSITION' in sem:
attribute = '[[position]]'
elif 'SV_RENDERTARGETARRAYINDEX' in sem:
attribute = '[[render_target_array_index]]'
shader_src += ['#line {}\n'.format(line_index)]
shader_src += [get_whitespace(line), dtype, ' ', name, ' ', attribute, ';\n']
continue
# metal requires tessellation factors of dtype half
if shader.stage == Stages.TESC or shader.stage == Stages.TESE:
dtype, name, sem = getMacro(line)
sem = sem.upper()
if sem == 'SV_TESSFACTOR' or sem == 'SV_INSIDETESSFACTOR':
line = line.replace(dtype, 'half')
# since we directly use the dtype of output patch for the mtl patch_control_point template,
# the inner struct needs attributes
if shader.stage == Stages.TESE:
if struct == shader.output_patch_arg[0]:
dtype, name, sem = getMacro(line)
elem = tuple(getMacro(line))
attribute_index = str(shader.structs[struct].index(elem))
line = line.replace(name, name + ' [[attribute(' + attribute_index + ')]]')
# handle cbuffer and pushconstant elements
if (parsing_cbuffer or parsing_pushconstant) and line.strip().startswith('DATA('):
dt, name, sem = getMacro(line)
element_basename = getArrayBaseName(name)
macro = get_uid(element_basename)
shader_src += ['#define ', macro, '\n']
if parsing_cbuffer:
elemen_path = parsing_cbuffer[0] + '.' + element_basename
is_embedded = parsing_cbuffer[1] in ab_elements
# for non-embedded cbuffer, access directly using struct access
global_reference_paths[element_basename] = parsing_cbuffer[0]
global_reference_args[element_basename] = 'constant struct ' + parsing_cbuffer[0] + '& ' + parsing_cbuffer[0]
if is_embedded:
elemen_path = 'AB_' + parsing_cbuffer[1] + '.' + elemen_path
global_reference_paths[element_basename] = 'AB_' + parsing_cbuffer[1]
global_reference_args[element_basename] = 'constant struct AB_' + parsing_cbuffer[1] + '& ' + 'AB_' + parsing_cbuffer[1]
if parsing_pushconstant:
elemen_path = parsing_pushconstant[0] + '.' + element_basename
global_reference_paths[element_basename] = parsing_pushconstant[0]
global_reference_args[element_basename] = 'constant struct ' + parsing_pushconstant[0] + '& ' + parsing_pushconstant[0]
shader_src += ['#define _Get_', element_basename, ' ', elemen_path, '\n']
if is_groupshared_decl(line):
dtype, dname = getMacro(line)
basename = getArrayBaseName(dname)
shader_src += ['\n\t// Metal GroupShared Declaration: ', basename, '\n']
macro = get_uid(basename)
shader_src += ['#define ', macro, '\n']
entry_declarations += [(macro, ['threadgroup {} {};'.format(dtype, dname)])]
array_decl = get_array_decl(dname)
global_reference_paths[basename] = basename
global_reference_args[basename] = 'threadgroup ' + dtype + ' (&' + basename + ')' + array_decl
shader_src += ['#define _Get_', basename, ' ', basename, '\n']
shader_src += ['#line {}\n'.format(line_index)]
shader_src += ['\t// End of GroupShared Declaration: ', basename, '\n']
continue
if 'PUSH_CONSTANT' in line:
parsing_pushconstant = tuple(getMacro(line))
struct_uid = get_uid(parsing_pushconstant[0])
shader_src += ['#define ', struct_uid, '\n']
if '};' in line and parsing_pushconstant:
shader_src += [line]
push_constant_decl = parsing_pushconstant
pushconstant_name = push_constant_decl[0]
push_constant_location = '[[buffer(UPDATE_FREQ_USER)]]'
mainArgs += [(struct_uid, ['constant struct ', pushconstant_name, '& ', pushconstant_name, ' ', push_constant_location])]
parsing_pushconstant = None
struct_references = []
struct_uid = None
last_res_decl = len(shader_src)+1
continue
if 'CBUFFER' in line:
fsl_assert(parsing_cbuffer == None, message='Inconsistent cbuffer declaration: \"' + line + '\"')
parsing_cbuffer = tuple(getMacro(line))
cbuffer_decl = parsing_cbuffer
struct_uid = get_uid(cbuffer_decl[0])
shader_src += ['#define ', struct_uid, '\n']
if '};' in line and parsing_cbuffer:
shader_src += [line]
cbuffer_decl = parsing_cbuffer
cbuffer_name, cbuffer_freq, dxreg = cbuffer_decl[:3]
is_rootcbv = 'rootcbv' in cbuffer_name
if cbuffer_freq not in metal_ab_frequencies and not is_rootcbv:
shader_src += ['#line {}\n'.format(line_index)]
shader_src += ['\t// Ignored CBuffer Declaration: '+line+'\n']
continue
is_embedded = cbuffer_freq in ab_elements
if not is_embedded:
location = buffer_location
if is_rootcbv:
location = 'UPDATE_FREQ_USER + 1'
else:
buffer_location += 1
mainArgs += [(struct_uid, ['constant struct ', cbuffer_name, '& ', cbuffer_name, ' [[buffer({})]]'.format(location)])]
else:
ab_elements[cbuffer_freq] += [(struct_uid, ['constant struct ', cbuffer_name, '& ', cbuffer_name])]
parsing_cbuffer = None
struct_references = []
struct_uid = None
last_res_decl = len(shader_src)+1
continue
# consume resources
if 'RES(' in line:
resource = tuple(getMacro(line))
resType, resName, freq, dxreg = resource[:4]
baseName = getArrayBaseName(resName)
macro = get_uid(baseName)
shader_src += ['#define ', macro, '\n']
is_embedded = freq in ab_elements and freq != 'UPDATE_FREQ_USER'
if freq not in metal_ab_frequencies:
shader_src += ['#line {}\n'.format(line_index)]
shader_src += ['\t// Ignored Resource Declaration: '+line+'\n']
continue
if not is_embedded: # regular resource
shader_src += ['\n\t// Metal Resource Declaration: ', line, '\n']
prefix = ''
postfix = ' '
ctor_arg = None
if 'Buffer' in resType:
prefix = 'device ' if 'RW' in resType else 'constant '
postfix = '* '
ctor_arg = ['', prefix, resType, '* ', resName]
else:
ctor_arg = ['thread ', resType, '& ', resName]
if 'Sampler' in resType:
binding = ' [[sampler({})]]'.format(sampler_location)
sampler_location += 1
elif 'Tex' in resType or 'Depth' in getMacroName(resType):
binding = ' [[texture({})]]'.format(texture_location)
texture_location += 1
elif 'Buffer' in resType:
binding = ' [[buffer({})]]'.format(buffer_location)
buffer_location += 1
else:
fsl_assert(False, message="Unknown Resource location")
main_arg = [prefix , resType, postfix, baseName, binding, ' // main arg']
mainArgs += [(macro, main_arg)]
global_reference_paths[baseName] = baseName
if not isArray(resName):
global_reference_args[baseName] = 'thread ' + resType + '& ' + resName
if 'Buffer' in resType:
space = 'constant'
if 'RW' in resType:
space = 'device'
global_reference_args[baseName] = space + ' ' + resType + '* ' + resName
else:
array = resName[resName.find('['):]
global_reference_args[baseName] = 'thread ' + resType + '(&' + baseName+') ' + array
shader_src += ['#define _Get_', baseName, ' ', baseName, '\n']
shader_src += ['\t// End of Resource Declaration: ', resName, '\n']
else: # resource is embedded in argbuf
shader_src += ['\n\t// Metal Embedded Resource Declaration: ', line, '\n']
argBufType = 'AB_' + freq
basename = getArrayBaseName(resName)
if 'Buffer' in resType:
space = ' device ' if resType.startswith('RW') else ' constant '
ab_element = [space, resType, '* ', resName]
else:
ab_element = [resType, ' ', resName]
ab_elements[freq] += [(macro, ab_element)]
global_reference_paths[baseName] = argBufType
if not isArray(resName):
global_reference_args[baseName] = 'constant ' + resType + '& ' + resName
if 'Buffer' in resType:
global_reference_args[baseName] = 'constant ' + resType + '* ' + resName
else:
array = resName[resName.find('['):]
global_reference_args[baseName] = 'constant ' + resType + '(&' + baseName+') ' + array
global_reference_args[baseName] = 'constant struct ' + argBufType + '& ' + argBufType
shader_src += ['#define _Get_', baseName, ' ', argBufType, '.', baseName, '\n']
shader_src += ['\t//End of Resource Declaration: ', baseName, '\n']
last_res_decl = len(shader_src)+1
shader_src += ['#line {}\n'.format(line_index)]
continue
# create comment hint for shader reflection
if 'NUM_THREADS(' in line:
elems = getMacro(line)
for i, elem in enumerate(elems):
if not elem.isnumeric():
assert elem in shader.defines, "arg {} to NUM_THREADS needs to be defined!".format(elem)
elems[i] = shader.defines[elem]
line = '// [numthreads({}, {}, {})]\n'.format(*elems)
# convert shader entry to member function
if '_MAIN(' in line:
parsing_main = len(shader_src)
ab_decl = declare_argument_buffers(mainArgs)
ab_decl_location = last_res_decl if not explicit_res_decl else explicit_res_decl
shader_src = shader_src[:ab_decl_location] + ab_decl + shader_src[ab_decl_location:]
if shader.stage == Stages.TESE:
num_control_points = shader.output_patch_arg[1]
patch_type = get_mtl_patch_type(tessellation_layout)
shader_src += ['[[patch(', patch_type, ', ', num_control_points, ')]]\n']
# insert VSMain code
if shader.stage == Stages.TESC:
shader_src += vs_main
shader_src += ['//[numthreads(32, 1, 1)]\n']
mtl_returntype = 'void'
if shader.returnType and shader.stage != Stages.TESC:
mtl_returntype = getMacroName(shader.returnType)
shader_src += [msl_target, ' ', mtl_returntype, ' stageMain(\n']
prefix = '\t'
if shader.stage == Stages.TESE:
if shader.output_patch_arg:
dtype, _, dname = shader.output_patch_arg
shader_src += [prefix+'patch_control_point<'+dtype+'> '+dname+'\n']
prefix = '\t,'
elif shader.stage == Stages.TESC:
shader_src += [prefix+'uint threadId [[thread_position_in_grid]]\n']
prefix = '\t,'
if shader.input_patch_arg:
dtype, _, dname = shader.input_patch_arg
else:
for dtype, var in shader.struct_args:
shader_src += [ prefix+dtype+' '+var+'[[stage_in]]\n']
prefix = '\t,'
for dtype, dvar in shader.flat_args:
if 'SV_OUTPUTCONTROLPOINTID' in dtype.upper(): continue
innertype = getMacro(dtype)
semtype = getMacroName(dtype)
shader_src += [prefix+innertype+' '+dvar+' '+semtype.upper()+'\n']
prefix = '\t,'
if shader.enable_waveops:
shader_src += [prefix, 'uint simd_lane_id [[thread_index_in_simdgroup]]\n']
prefix = '\t,'
for macro, arg in mainArgs:
shader_src += ['#ifdef ', macro, '\n']
shader_src += [prefix, *arg, '\n']
prefix = '\t,'
shader_src += ['#endif\n']
shader_src += [')\n']
shader_src += ['#line {}\n'.format(line_index)]
continue
if 'INIT_MAIN' in line:
for macro, entry_declaration in entry_declarations:
shader_src += ['#ifdef ', macro,'\n']
shader_src += ['\t', *entry_declaration, '\n']
shader_src += ['#endif\n']
if shader.stage == Stages.TESC:
# fetch VS data, call VS main
dtype, dim, dname = shader.input_patch_arg
shader_src += [
'\n\tif (threadId > drawInfo->vertexCount) return;\n',
'\n\t// call VS main\n',
'\tconst '+dtype+' '+dname+'['+dim+'] = { VSMain(vertexInput[threadId]) };\n',
]
for dtype, dvar in shader.flat_args:
if 'SV_OUTPUTCONTROLPOINTID' in dtype.upper():
shader_src += ['\t'+getMacro(dtype)+' '+getMacro(dvar)+' = 0u;\n'] # TODO: extend this for >1 CPs
shader_src += ['#line {}\n'.format(line_index), '//'+line]
continue
if re.search('(^|\s+)RETURN', line):
ws = get_whitespace(line)
return_statement = [ ws+'{\n' ]
# void entry, return nothing
if not shader.returnType:
return_statement += [ws+'\treturn;\n']
else:
return_value = getMacro(line)
# for tessellation, write tesc results to output buffer and call PCF
if shader.stage == Stages.TESC:
return_statement += [ws+'\thullOutputBuffer[threadId] = '+return_value+';\n']
return_statement += [ws+'\ttessellationFactorBuffer[threadId] = '+shader.pcf+'(\n']
prefix = ws+'\t\t'
for dtype, dvar in shader.pcf_arguments:
if 'INPUT_PATCH' in dtype:
return_statement += [prefix+shader.input_patch_arg[-1]+'\n']
prefix = ws+'\t\t,'
if 'SV_PRIMITIVEID' in dtype.upper():
return_statement += [prefix+'0u'+'\n'] # TODO: extend this for >1 CPs
prefix = ws+'\t\t,'
return_statement += [ws+'\t);\n']
# entry declared with returntype, return var
else:
return_statement += [ws+'\treturn '+return_value+';\n']
return_statement += [ ws+'}\n' ]
shader_src += return_statement
shader_src += ['#line {}\n'.format(line_index), '//'+line]
continue
# tessellation PCF
if shader.pcf and ' '+shader.pcf in line:
for dtype, dvar in shader.pcf_arguments:
if 'INPUT_PATCH' in dtype: continue
innertype = getMacro(dtype)
line = line.replace(dtype, innertype)
# since we modify the pcf signature, need to update the global_fn_table accordingly
_, (line_no, _) = global_fn_table[shader.pcf]
global_fn_table[shader.pcf] = (line, (line_no, line.find(shader.pcf)+len(shader.pcf)+1))
if shader.pcf and 'PCF_INIT' in line:
line = get_whitespace(line)+'//'+line.strip()+'\n'
if shader.pcf and 'PCF_RETURN' in line:
ws = get_whitespace(line)
return_value = getMacro(line)
line = ws+'return '+return_value+';\n'
if shader_src_len != len(shader_src):
shader_src += ['#line {}\n'.format(line_index)]
shader_src += [line]
shader_src += ['\n']
# for each function, expand signature and calls to pass global resource references
for fn, references in global_references.items():
insert_line, (_, insert_loc) = global_fn_table[fn]
call_additions = []
signature_additions = []
for reference in references:
if global_reference_paths[reference] not in call_additions:
call_additions += [global_reference_paths[reference]]
signature_additions += [global_reference_args[reference]]
modified_signature = False
for i, line in enumerate(shader_src):
if line.strip().startswith('//'): continue
# modify signatures
l_call = line.find(fn+'(')
if insert_line in line:
for parameter in signature_additions:
if line[insert_loc-1:insert_loc+1] == '()':
line = line[:insert_loc] + parameter + line[insert_loc:]
else:
line = line[:insert_loc] + parameter + ', ' + line[insert_loc:]
# print('modify signature:', fn, shader_src[i], line)
shader_src[i] = line
modified_signature = True
# modify calls
elif modified_signature and l_call > 0 and line[l_call-1] in ' =\t(!':
l2 = line.find(');', l_call)
l2 = 0
counter = 0
for j, c in enumerate(line[l_call+len(fn):]):
if c == '(':
counter+=1
if counter == 1:
l2 = j+l_call+len(fn)+1
break
if c == ')':
counter-=1
for argument in call_additions:
if line[l2-1:l2+1] == '()':
line = line[:l2] + argument + line[l2:]
else:
line = line[:l2] + argument + ', ' + line[l2:]
# print('modify call:', shader_src[i], line)
shader_src[i] = line
open(dst, 'w').writelines(shader_src)
# sys.exit(0)
return 0
def gles(fsl, dst):
shader = getShader(fsl, dst, line_directives=False)
#Only vertex and fragment shaders are valid for OpenGL ES 2.0
if shader.stage != Stages.VERT and shader.stage != Stages.FRAG:
print(
"Invalid OpenGL ES 2.0 shader given, only .vert and .frag shaders are valid."
)
return 1
shader_src = getHeader(fsl)
shader_src += [
'#version 100\n', '\nprecision highp float;\nprecision highp int;\n\n'
]
shader_src += ['#define STAGE_', shader.stage.name, '\n']
shader_src += ['#define GLES\n']
header_path = os.path.join(os.path.dirname(os.path.dirname(__file__)),
'includes', 'gles.h')
header_lines = open(header_path).readlines()
shader_src += header_lines + ['\n']
if True or not os.path.exists(dst):
os.makedirs(os.path.dirname(dst), exist_ok=True)
# retrieve output patch size
patch_size = 0
for line in shader.lines:
if 'OUTPUT_CONTROL_POINTS' in line:
patch_size = int(getMacro(line))
arrBuffs = [
'Get(' + getArrayBaseName(res[1]) + ')' for res in shader.resources
if 'Buffer' in res[0] and (isArray(res[1]))
]
returnType = None if not shader.returnType else (getMacroName(
shader.returnType), getMacro(shader.returnType))
skip_line = False
parsing_struct = None
parsing_comments = False
parsing_ubuffer = None
parsing_ubuffer_float_count = 0
nonuniformresourceindex = None
parsed_entry = False
struct_construction = {}
get_declarations = {}
struct_declarations = []
input_assignments = []
return_assignments = []
def is_struct(var_type):
return var_type in struct_construction
def set_ubo(ubo_name, basename, elemname, elem_type, getname, float_stride,
fromStruct, isArray, float_offset, result):
if (is_struct(elem_type)):
structItem = struct_construction[elem_type]
for uniformIndex in range(structItem['uniformCount']):
elem_dtype, elem_name, _ = getMacro(structItem[uniformIndex])
struct_get_name = getname + '_' + elem_name.upper()
float_offset, result = set_ubo(
ubo_name, basename,
elemname + ('\.' if isArray else '.') + elem_name,
elem_dtype, struct_get_name, float_stride, True, isArray,
float_offset, result) #recurse
else:
if isArray:
if fromStruct:
get_declarations['Get\(' + basename + '\)\[.+\]' +
elemname] = (getname + '(#)', True)
else:
get_declarations['Get\(' + basename +
'\)\[.+\]'] = (getname + '(#)', True)
else:
if fromStruct:
get_declarations['Get(' + basename + ')' +
elemname] = (getname, False)
else:
get_declarations['Get(' + basename + ')'] = (getname,
False)
elem_float_size = getFloatSize(elem_type, struct_construction)
element_path = setUBOConversion(elem_type, ubo_name, float_offset,
float_stride, isArray,
struct_construction)
if isArray:
replaced_value = element_path.replace('#', 'X')
result += '#define ' + getname + '(X) ' + replaced_value + '\n'
else:
result += '#define ' + getname + ' ' + element_path + '\n'
float_offset += elem_float_size
return float_offset, result
def declare_buffer(fsl_declaration):
#Set a default max buffer size, GLSL v100 does not allow unknown array sizes
#TODO: We might want to use a texture as buffer instead to allow unknown large buffer sizes (requires engine changes)
default_max_buffer_size = 256
buffer_type, name, _, _, _ = fsl_declaration
data_type = getMacro(buffer_type)
assert not isArray(name), 'Cannot use array of buffers in glsl v100'
result = []
stride = getFloatSize(data_type, struct_construction)
get_name = 'Get_' + name.upper()
_, result = set_ubo(name, name, '', data_type, get_name, stride, False,
True, 0, result)
result += 'uniform float4 {}[{}];\n'.format(
name, str(default_max_buffer_size))
return result
for line in shader.lines:
def get_uid(name):
return '_' + name + '_' + str(len(shader_src))
# dont process commented lines
if line.strip().startswith('//'):
shader_src += [line]
continue
if line.strip().startswith('/*'):
parsing_comments = True
continue
if line.strip().startswith('*/'):
parsing_comments = False
continue
if parsing_comments:
continue
if 'INDIRECT_DRAW()' in line:
continue
# TODO: handle this differently
if '#ifdef NO_FSL_DEFINITIONS' in line:
skip_line = True
if skip_line and '#endif' in line:
skip_line = False
continue
if skip_line:
continue
if line.strip().startswith('STRUCT('):
parsing_struct = getMacro(line)
struct_construction[parsing_struct] = {
'floatSize': 0,
'uniformCount': 0
}
if parsing_struct and '};' in line:
if not line.endswith('\n'): line += '\n'
shader_src += [line]
print("{} struct size = {}".format(
parsing_struct,
str(struct_construction[parsing_struct]['floatSize'])))
for macro, struct_declaration in struct_declarations:
shader_src += ['#ifdef ', macro, '\n']
shader_src += [*struct_declaration, '\n']
shader_src += ['#endif\n']
struct_declarations = []
parsing_struct = None
continue
# handle struct data declarations
if parsing_struct and line.strip().startswith('DATA('):
# handle entry return type
if shader.returnType and parsing_struct in shader.returnType:
var = 'out_' + shader.returnType
elem_dtype, elem_name, sem = getMacro(line)
sem = sem.upper()
basename = getArrayBaseName(elem_name)
macro = get_uid(basename)
shader_src += ['#define ', macro, '\n']
output_datapath = elem_name
reference = None
if sem == 'SV_POSITION':
output_datapath = 'gl_Position'
elif sem == 'SV_POINTSIZE':
output_datapath = 'gl_PointSize'
elif sem == 'SV_DEPTH':
output_datapath = 'gl_FragDepth'
else:
reference = [
'RES_OUT(', elem_dtype, ', ', output_datapath, ');'
]
# unroll attribute arrays
if isArray(elem_name):
assignment = [
var, '_', basename, ' = ', var, '.', basename, ''
]
else:
assignment = [output_datapath, ' = ', var, '.', elem_name]
if reference:
struct_declarations += [(macro, reference)]
return_assignments += [(macro, assignment)]
# elif shader.entry_arg and parsing_struct in shader.entry_arg:
elif is_input_struct(parsing_struct, shader):
var = get_input_struct_var(parsing_struct, shader)
elem_dtype, elem_name, sem = getMacro(line)
sem = sem.upper()
flat_modifier = 'FLAT(' in line
if flat_modifier:
elem_dtype = getMacro(elem_dtype)
line = get_whitespace(line) + getMacro(
elem_dtype) + ' ' + elem_name + ';\n'
noperspective_modifier = 'noperspective(' in line
if noperspective_modifier:
elem_dtype = getMacro(elem_dtype)
line = get_whitespace(line) + getMacro(
elem_dtype) + ' ' + elem_name + ';\n'
is_array = isArray(elem_name)
basename = getArrayBaseName(elem_name)
macro = get_uid(basename)
shader_src += ['#define ', macro, '\n']
input_datapath = var + '_' + elem_name
if sem == 'SV_POINTSIZE': continue
input_value = None
if shader.stage == Stages.VERT:
# for vertex shaders, use the semantic as attribute name (for name-based reflection)
input_value = getSemanticName(sem)
if not input_value:
input_value = elem_name
reference = [
'RES_IN(', elem_dtype, ', ', input_value, ');'
]
elif shader.stage == Stages.FRAG:
# for fragment shaders, set the input as RES_OUT to be a varying
input_value = elem_name
reference = [
'RES_OUT(', elem_dtype, ', ', input_value, ');'
]
if flat_modifier:
reference.insert(0, 'flat ')
if noperspective_modifier:
reference.insert(0, 'noperspective ')
# input semantics
if sem == 'SV_POSITION' and shader.stage == Stages.FRAG:
input_value = elem_dtype + '(float4(gl_FragCoord.xyz, 1.0 / gl_FragCoord.w))'
reference = []
assignment = []
if sem == 'SV_VERTEXID': input_value = 'gl_VertexIndex'
# unroll attribute arrays
# if is_array:
# assignment = ['for(int i=0; i<', str(getArrayLenFlat(elem_name)), ';i++) ']
# assignment += [var, '.', basename, '[i] = ', var, '_', basename, '[i]']
# else:
assignment = [var, '.', basename, ' = ', input_value]
if reference:
struct_declarations += [(macro, reference)]
input_assignments += [(macro, assignment)]
else:
# Store struct information for later usage in declaring UBOs
elem_dtype, elem_name, sem = getMacro(line)
struct_construction[parsing_struct][
'floatSize'] += getFloatSize(elem_dtype,
struct_construction)
uniformCount = struct_construction[parsing_struct][
'uniformCount']
struct_construction[parsing_struct][uniformCount] = line
struct_construction[parsing_struct]['uniformCount'] += 1
# Handle uniform buffers
if parsing_ubuffer and line.strip().startswith('DATA('):
elem_dtype, name, sem = getMacro(line)
element_basename = getArrayBaseName(name)
result = []
float_stride = getFloatSize(elem_dtype, struct_construction)
array_length = 1
if isArray(name):
array_length = getArrayLen(shader.defines, name)
get_name = 'Get_' + element_basename.upper()
_, result = set_ubo(parsing_ubuffer, element_basename, '',
elem_dtype, get_name, float_stride, False,
isArray(name), parsing_ubuffer_float_count,
result)
parsing_ubuffer_float_count += float_stride * (array_length)
shader_src += result
continue
if 'CBUFFER' in line:
fsl_assert(parsing_ubuffer == None,
message='Inconsistent cbuffer declaration: \"' + line +
'\"')
parsing_ubuffer, _, _, _ = getMacro(line)
continue
if 'PUSH_CONSTANT' in line:
fsl_assert(parsing_ubuffer == None,
message='Inconsistent push_constant declaration: \"' +
line + '\"')
parsing_ubuffer, _ = getMacro(line)
continue
if '{' in line and parsing_ubuffer:
continue
if '};' in line and parsing_ubuffer:
parsing_ubuffer_float_count = math.ceil(
parsing_ubuffer_float_count / 4)
shader_src += [
'uniform float4 ', parsing_ubuffer, '[',
str(parsing_ubuffer_float_count), '];\n'
]
parsing_ubuffer_float_count = 0
parsing_ubuffer = None
continue
resource_decl = None
if line.strip().startswith('RES('):
resource_decl = getMacro(line)
print(resource_decl)
fsl_assert(len(resource_decl) == 5,
fsl,
message='invalid Res declaration: \'' + line + '\'')
basename = getArrayBaseName(resource_decl[1])
get_name = 'Get_' + basename.upper()
# No samplers available in GLSL v100, define NO_SAMPLER
if resource_decl and is_sampler(resource_decl):
get_declarations['Get(' + basename + ')'] = (get_name, False)
shader_src += ['#define ' + get_name + ' NO_SAMPLER' '\n']
continue
# Handle buffer resource declarations for GLES
if is_buffer(resource_decl):
shader_src += declare_buffer(resource_decl)
continue
basename = getArrayBaseName(resource_decl[1])
get_name = 'Get_' + basename.upper()
get_declarations['Get(' + basename + ')'] = (get_name, False)
shader_src += ['#define ' + get_name + ' ' + basename + '\n']
#TODO handle references to arrays of structured buffers for GLES
#for buffer_array in arrBuffs:
# line = insert_buffer_array_indirections(line, buffer_array)
#TODO specify format qualified image resource for GLES
if resource_decl and is_rw_texture(resource_decl):
#shader_src += declare_rw_texture(resource_decl)
continue
#TODO Check if usage is needed within GLES
if 'BeginNonUniformResourceIndex(' in line:
#index, max_index = getMacro(line), None
#assert index != [], 'No index provided for {}'.format(line)
#if type(index) == list:
# max_index = index[1]
# index = index[0]
#nonuniformresourceindex = index
## if type(max_index) is str:
#if max_index and not max_index.isnumeric():
# max_index = ''.join(c for c in shader.defines[max_index] if c.isdigit())
#shader_src += BeginNonUniformResourceIndex(nonuniformresourceindex, max_index)
continue
if 'EndNonUniformResourceIndex()' in line:
#assert nonuniformresourceindex, 'EndNonUniformResourceIndex: BeginNonUniformResourceIndex not called/found'
#shader_src += EndNonUniformResourceIndex(nonuniformresourceindex)
#nonuniformresourceindex = None
continue
if nonuniformresourceindex:
#shader_src += [line[:-1], ' \\\n']
continue
if '_MAIN(' in line:
for dtype, dvar in shader.flat_args:
sem = getMacroName(dtype).upper()
if sem == 'SV_INSTANCEID':
shader_src += ['uniform int ', sem, ';\n\n']
if shader.returnType and shader.stage == Stages.VERT and shader.returnType not in shader.structs:
shader_src += [
'RES_OUT(', shader.returnType, ', ', 'out_',
shader.returnType, ');\n'
]
#TODO Check if this is needed somewere
#if shader.input_patch_arg:
#patch_size = shader.input_patch_arg[1]
#shader_src += ['layout(vertices = ', patch_size, ') out;\n']
shader_src += ['void main()\n']
parsed_entry = True
continue
if parsed_entry and re.search('(^|\s+)RETURN', line):
ws = get_whitespace(line)
output_statement = [ws + '{\n']
if shader.returnType:
output_value = getMacro(line)
if shader.returnType not in shader.structs:
if shader.stage == Stages.FRAG:
output_statement += [
ws + '\tgl_FragColor = ' + output_value + ';\n'
]
else:
output_statement += [
ws + '\tout_' + shader.returnType + ' = ' +
output_value + ';\n'
]
else:
output_statement += [
ws + '\t' + shader.returnType + ' out_' +
shader.returnType + ' = ' + output_value + ';\n'
]
for macro, assignment in return_assignments:
output_statement += ['#ifdef ', macro, '\n']
output_statement += [ws + '\t', *assignment, ';\n']
output_statement += ['#endif\n']
output_statement += [ws + '}\n']
shader_src += output_statement
continue
if 'INIT_MAIN' in line:
# assemble input
for dtype, var in shader.struct_args:
if shader.input_patch_arg and dtype in shader.input_patch_arg[
0]:
dtype, dim, var = shader.input_patch_arg
shader_src += ['\t', dtype, ' ', var, '[', dim, '];\n']
continue
if shader.output_patch_arg and dtype in shader.output_patch_arg[
0]:
dtype, dim, var = shader.output_patch_arg
shader_src += ['\t', dtype, ' ', var, '[', dim, '];\n']
continue
shader_src += ['\t', dtype, ' ', var, ';\n']
for macro, assignment in input_assignments:
shader_src += ['#ifdef ', macro, '\n']
shader_src += ['\t', *assignment, ';\n']
shader_src += ['#endif\n']
''' additional inputs '''
for dtype, dvar in shader.flat_args:
innertype = getMacro(dtype)
semtype = getMacroName(dtype)
shader_src += [
'\t' + innertype + ' ' + dvar + ' = ' + innertype + '(' +
semtype.upper() + ');\n'
]
''' generate a statement for each vertex attribute
this should not be necessary, but it influences
the attribute order for the SpirVTools shader reflection
'''
# if shader.stage == Stages.VERT:# and shader.entry_arg:
# for dtype, var in shader.struct_args:
# for _, n, s in shader.structs[dtype]:
# if s.upper() == 'SV_VERTEXID': continue
# shader_src += ['\t', s ,';\n']
continue
updatedline = line
for key, value in get_declarations.items():
if value[1] and isArray(updatedline):
# Regex match and replace
replace_value = value[0].replace('#',
getArrayLenFlat(updatedline))
updatedline = re.sub(key, replace_value, updatedline)
elif key in updatedline:
updatedline = updatedline.replace(key, value[0])
#Strip .f -> .0 | 0f -> 0 float value definitions, those are not supported on GLES
def replacef(matchReg):
if len(matchReg.group(0)) > 2:
return matchReg.group(0)[:2] + '0'
return matchReg.group(0)[0]
shader_src += [re.sub('\d\.?f', replacef, updatedline)]
open(dst, 'w').writelines(shader_src)
return 0