blob: 9799b0a36cf0572089a0ec99d98c7e0109287a5a [file] [log] [blame]
{{- /*
--------------------------------------------------------------------------------
Template file for use with tools/builtin-gen to generate the wgsl files in the
./literal/... and ./var/... subdirectories
See:
* tools/cmd/intrinsic-gen/gen for structures used by this template
* https://golang.org/pkg/text/template/ for documentation on the template syntax
--------------------------------------------------------------------------------
*/ -}}
{{- /* For each permutation of each overload of each function... */ -}}
{{- range .Sem.Builtins -}}
{{- /* TODO(crbug.com/tint/1536): Remove the bodge below after we support [[extension("extension_name")]] in intrinsics.def */}}
{{- if not (or (eq .Name "dot4I8Packed") (eq .Name "dot4U8Packed"))}}
{{- range .Overloads -}}
{{- range Permute . -}}
{{- /* Generate a ./literal/<function>/<permuataion-hash>.wgsl file using
the Permutation macro defined below */ -}}
{{- $file := printf "./literal/%v/%v.wgsl" .Intrinsic.Name .Hash -}}
{{- $content := Eval "Permutation" "Overload" . "Mode" "literal" -}}
{{- WriteFile $file $content -}}
{{- /* Generate a ./var/<function>/<permuataion-hash>.wgsl file using
the Permutation macro defined below */ -}}
{{- $file := printf "./var/%v/%v.wgsl" .Intrinsic.Name .Hash -}}
{{- $content := Eval "Permutation" "Overload" . "Mode" "var" -}}
{{- WriteFile $file $content -}}
{{- end }}
{{- end }}
{{- end }}
{{- end }}
{{- /* ------------------------------------------------------------------ */ -}}
{{- define "Permutation" -}}
{{- /* Emits the body of the intrinsic permuation .wgsl file */ -}}
{{- /* ------------------------------------------------------------------ */ -}}
{{- $overload := .Overload -}}
{{- $mode := .Mode -}}
{{- $builtin := $overload.Intrinsic.Name -}}
{{- $permutation := printf "%v_%v" $builtin $overload.Hash -}}
{{- $args := Map -}}
{{- /* Generate RW storage buffer parameters */ -}}
{{- $sb_rw_fields := Eval "EmitBufferFields" "overload" $overload
"var_name" "sb_rw"
"storage" "storage"
"access" "read_write"
"args" $args -}}
{{- if $sb_rw_fields -}}
struct SB_RW {
{{- $sb_rw_fields -}}
};
@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
{{ end -}}
{{- /* Generate RO storage buffer parameters */ -}}
{{- $sb_ro_fields := Eval "EmitBufferFields" "overload" $overload
"var_name" "sb_ro"
"storage" "storage"
"access" "read"
"args" $args -}}
{{- if $sb_ro_fields -}}
struct SB_RO {
{{- $sb_ro_fields -}}
};
@group(0) @binding(1) var<storage, read> sb_ro : SB_RO;
{{ end -}}
{{- /* Generate uniform buffer parameters */ -}}
{{- $ub_fields := Eval "EmitBufferFields" "overload" $overload
"var_name" "ub"
"storage" "uniform"
"access" "read"
"args" $args -}}
{{- if $ub_fields -}}
struct UB {
{{- $ub_fields -}}
};
@group(0) @binding(1) var<uniform> ub : UB;
{{ end -}}
{{- /* Generate module-scoped handle variables */ -}}
{{- range $i, $p := $overload.Parameters }}
{{- $class := Eval "StorageClass" $p.Type -}}
{{- if eq "ptr" $p.Type.Target.Name -}}
{{- $el_type := Eval "Type" (index $p.Type.TemplateArguments 1)}}
{{- if eq "handle" $class -}}
@group(1) @binding({{$i}}) var arg_{{$i}}: {{$el_type}};
{{ $args.Put $i (printf "&arg_%v" $i) -}}
{{- else if eq "workgroup" $class -}}
var<workgroup> arg_{{$i}}: {{$el_type}};
{{ $args.Put $i (printf "&arg_%v" $i) -}}
{{- else if eq "private" $class -}}
var<private> arg_{{$i}}: {{$el_type}};
{{ $args.Put $i (printf "&arg_%v" $i) -}}
{{- end -}}
{{- else -}}
{{- $type := Eval "Type" $p.Type}}
{{- if eq "handle" $class -}}
@group(1) @binding({{$i}}) var arg_{{$i}}: {{$type}};
{{ $args.Put $i (printf "arg_%v" $i) -}}
{{- else if eq "workgroup" $class -}}
var<workgroup> arg_{{$i}}: {{$type}};
{{ $args.Put $i (printf "arg_%v" $i) -}}
{{- else if eq "private" $class -}}
var<private> arg_{{$i}}: {{$type}};
{{ $args.Put $i (printf "arg_%v" $i) -}}
{{- end -}}
{{- end -}}
{{- end -}}
{{- /* Generate the function that calls the intrinsic */ -}}
{{- /*newline*/}}
// {{$.Overload}}
fn {{$permutation}}() {
{{/* Build the parameters either as 'var' or inline values */ -}}
{{- range $i, $p := $overload.Parameters -}}
{{- $class := Eval "StorageClass" $p.Type -}}
{{- if eq "function" $class -}}
{{- if eq "ptr" $p.Type.Target.Name -}}
{{- /*indent*/}} var arg_{{$i}}: {{template "Type" index $p.Type.TemplateArguments 1}};
{{ $args.Put $i (printf "&arg_%v" $i) -}}
{{- else if and (not $p.IsConst) (eq "var" $mode) -}}
{{- /*indent*/}} var arg_{{$i}} = {{Eval "ArgumentValue" $p}};
{{ $args.Put $i (printf "arg_%v" $i) -}}
{{- else -}}
{{- $args.Put $i (Eval "ArgumentValue" $p) -}}
{{- end -}}
{{- end -}}
{{- end -}}
{{- /* Make the call to the intrinsic */ -}}
{{- /*indent*/}} {{/*indent*/ -}}
{{- if $overload.ReturnType -}}
var res{{if IsDeclarable $overload.ReturnType}}: {{template "Type" $overload.ReturnType}}{{end}} = {{/* preserve space after = */ -}}
{{- end -}}
{{$builtin}}(
{{- range $i, $p := $overload.Parameters -}}
{{- if $i -}}, {{end}}{{$args.Get $i -}}
{{- end -}}
);
}
{{/*new line*/ -}}
{{- if $overload.CanBeUsedInStage.Vertex }}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
{{$permutation}}();
return vec4<f32>();
}
{{ end -}}
{{- if $overload.CanBeUsedInStage.Fragment }}
@fragment
fn fragment_main() {
{{$permutation}}();
}
{{ end -}}
{{- if $overload.CanBeUsedInStage.Compute }}
@compute @workgroup_size(1)
fn compute_main() {
{{$permutation}}();
}
{{ end -}}
{{- end -}}
{{- /* ------------------------------------------------------------------ */ -}}
{{- define "EmitBufferFields" -}}
{{- /* Emits a struct with the fields that match the given storage class */ -}}
{{- /* and access. */ -}}
{{- /* Argument is a map with the following expected keys: */ -}}
{{- /* 'overload' - the current overload */ -}}
{{- /* 'var_name' - name of the variable of the structure type */ -}}
{{- /* 'storage' - filtered storage class */ -}}
{{- /* 'access' - filtered access */ -}}
{{- /* 'args' - argument map that's populated with the fields */ -}}
{{- /* ------------------------------------------------------------------ */ -}}
{{- $overload := .Get "overload" -}}
{{- $var_name := .Get "var_name" -}}
{{- $filter_storage := .Get "storage" -}}
{{- $filter_access := .Get "access" -}}
{{- $args := .Get "args" -}}
{{- range $i, $p := $overload.Parameters }}
{{- $storage := Eval "StorageClass" $p.Type -}}
{{- $access := Eval "Access" $p.Type -}}
{{- if and (eq $filter_storage $storage) (eq $filter_access $access) }}
{{- if eq "ptr" $p.Type.Target.Name }}
arg_{{$i}}: {{template "Type" (index $p.Type.TemplateArguments 1)}},
{{ $args.Put $i (printf "&%v.arg_%v" $var_name $i) -}}
{{- else }}
arg_{{$i}}: {{template "Type" $p.Type}},
{{ $args.Put $i (printf "%v.arg_%v" $var_name $i) -}}
{{- end -}}
{{- end -}}
{{- end -}}
{{ end -}}
{{- /* ------------------------------------------------------------------ */ -}}
{{- define "StorageClass" -}}
{{- /* Returns the storage class for the given Fully Qualified Name */ -}}
{{- /* ------------------------------------------------------------------ */ -}}
{{- $name := .Target.Name -}}
{{- if eq $name "array" -}}storage
{{- else if HasPrefix $name "texture" -}}handle
{{- else if HasPrefix $name "sampler" -}}handle
{{- else if eq $name "ptr" -}}{{(index .TemplateArguments 0).Target.Name}}
{{- else -}}function
{{- end -}}
{{- end -}}
{{- /* ------------------------------------------------------------------ */ -}}
{{- define "Access" -}}
{{- /* Returns the access for the given Fully Qualified Name */ -}}
{{- /* ------------------------------------------------------------------ */ -}}
{{- $name := .Target.Name -}}
{{- if eq $name "ptr" -}}{{(index .TemplateArguments 2).Target.Name}}
{{- else -}}
{{- /* Emit the default for the storage class */ -}}
{{- /* https://gpuweb.github.io/gpuweb/wgsl/#storage-class */ -}}
{{- $storage := Eval "StorageClass" . -}}
{{- if eq $storage "function" -}}read_write
{{- else if eq $storage "private" -}}read_write
{{- else if eq $storage "workgroup" -}}read_write
{{- else if eq $storage "uniform" -}}read
{{- else if eq $storage "storage" -}}read
{{- else if eq $storage "handle" -}}read
{{- end -}}
{{- end -}}
{{- end -}}
{{- /* ------------------------------------------------------------------ */ -}}
{{- define "ArgumentValue" -}}
{{- /* Returns a value that can be used for the parameter argument */ -}}
{{- /* ------------------------------------------------------------------ */ -}}
{{- $ty := .Type -}}
{{- if eq $ty.Target.Name "i32" -}}
{{- /* If the parameter has the name 'level', then use '0' as the value. */ -}}
{{- /* Some texture arguments require the level parameter to be 0, and */ -}}
{{- /* constraint is not described in the definition file. */ -}}
{{- if eq .Name "level" -}}0
{{- else -}}1
{{- end -}}
{{- else if eq $ty.Target.Name "u32" -}}1u
{{- else if eq $ty.Target.Name "f32" -}}1.0
{{- else -}}{{template "Type" $ty}}()
{{- end -}}
{{- end -}}
{{- /* ------------------------------------------------------------------ */ -}}
{{- define "Type" -}}
{{- /* Emits the WGSL for the Fully Qualified Name argument */ -}}
{{- /* ------------------------------------------------------------------ */ -}}
{{- if IsType .Target -}}
{{- if eq .Target.Name "vec" -}}vec{{index .TemplateArguments 0}}<{{template "Type" index .TemplateArguments 1}}>
{{- else if eq .Target.Name "mat" -}}mat{{index .TemplateArguments 0}}x{{index .TemplateArguments 1}}<{{template "Type" index .TemplateArguments 2}}>
{{- else -}}{{.Target.Name}}{{template "TemplateArguments" .TemplateArguments}}
{{- end -}}
{{- else if IsEnumEntry .Target -}}{{.Target.Name}}
{{- else if IsEnumMatcher .Target -}}{{(index .Target.Options 0).Name}}
{{- else -}}<unhandled-fully-qualified-name-target={{- printf "%T" .Target -}}>
{{- end -}}
{{- end -}}
{{- /* ------------------------------------------------------------------ */ -}}
{{- define "TemplateArguments" -}}
{{- /* Emits the WGSL for the template argument list */ -}}
{{- /* ------------------------------------------------------------------ */ -}}
{{- if . -}}
<
{{- range $i, $a := . -}}
{{- if $i -}}, {{ end -}}
{{- if IsInt $a -}}{{- . -}}
{{- else -}}{{- template "Type" $a -}}
{{- end -}}
{{- end -}}
>
{{- end -}}
{{- end -}}