mirror of
https://github.com/encounter/dawn-cmake.git
synced 2025-08-03 10:45:36 +00:00
e6e96def66
Adapt the builtin parsing and resolving to also support operators. Will be used to generate intrinsic table entries for operators. This will simplify maintenance of the operators, and will greatly simplify the [AbstractInt -> i32|u32] [AbstractFloat -> f32|f16] logic. Bug: tint:1504 Change-Id: Id75735ea24e501877418812185796f3fba88a521 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/89026 Commit-Queue: Ben Clayton <bclayton@chromium.org> Kokoro: Kokoro <noreply+kokoro@google.com> Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
277 lines
12 KiB
Cheetah
277 lines
12 KiB
Cheetah
{{- /*
|
|
--------------------------------------------------------------------------------
|
|
Template file for use with tools/builtin-gen to generate the wgsl files in the
|
|
./gen/... 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 builtin... */ -}}
|
|
{{- range .Sem.Builtins -}}
|
|
{{- /* TODO(crbug.com/tint/1483): Remove the bodge below after smoothStep is removed from builtins.def */}}
|
|
{{- if not (eq .Name "smoothStep") }}
|
|
{{- range .Overloads -}}
|
|
{{- range Permute . -}}
|
|
{{- /* Generate a ./gen/<function>/<permuataion-hash>.wgsl file using
|
|
the Permutation macro defined below */ -}}
|
|
{{- $file := printf "./gen/%v/%v.wgsl" .Intrinsic.Name .Hash -}}
|
|
{{- $content := Eval "Permutation" . -}}
|
|
{{- WriteFile $file $content -}}
|
|
{{- end }}
|
|
{{- end }}
|
|
{{- end }}
|
|
{{- end }}
|
|
|
|
|
|
{{- /* ------------------------------------------------------------------ */ -}}
|
|
{{- define "Permutation" -}}
|
|
{{- /* Emits the body of the intrinsic permuation .wgsl file */ -}}
|
|
{{- /* ------------------------------------------------------------------ */ -}}
|
|
{{- $builtin := .Intrinsic.Name -}}
|
|
{{- $permutation := printf "%v_%v" $builtin .Hash -}}
|
|
{{- $args := Map -}}
|
|
|
|
{{- /* Generate RW storage buffer parameters */ -}}
|
|
{{- $sb_rw_fields := Eval "EmitBufferFields" "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" .
|
|
"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" .
|
|
"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 := .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 := .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 -}}
|
|
{{- $args.Put $i (Eval "ArgumentValue" $p) -}}
|
|
{{- end -}}
|
|
{{- end -}}
|
|
{{- end -}}
|
|
|
|
{{- /* Make the call to the intrinsic */ -}}
|
|
{{- /*indent*/}} {{/*indent*/ -}}
|
|
{{- if .ReturnType -}}
|
|
var res{{if IsDeclarable .ReturnType}}: {{template "Type" .ReturnType}}{{end}} = {{/* preserve space after = */ -}}
|
|
{{- end -}}
|
|
{{$builtin}}(
|
|
{{- range $i, $p := .Parameters -}}
|
|
{{- if $i -}}, {{end}}{{$args.Get $i -}}
|
|
{{- end -}}
|
|
);
|
|
}
|
|
{{/*new line*/ -}}
|
|
|
|
{{- if .CanBeUsedInStage.Vertex }}
|
|
@stage(vertex)
|
|
fn vertex_main() -> @builtin(position) vec4<f32> {
|
|
{{$permutation}}();
|
|
return vec4<f32>();
|
|
}
|
|
{{ end -}}
|
|
|
|
{{- if .CanBeUsedInStage.Fragment }}
|
|
@stage(fragment)
|
|
fn fragment_main() {
|
|
{{$permutation}}();
|
|
}
|
|
{{ end -}}
|
|
|
|
{{- if .CanBeUsedInStage.Compute }}
|
|
@stage(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 -}}
|