| {{- /* |
| -------------------------------------------------------------------------------- |
| Template file for use with tools/src/cmd/gen to generate the wgsl files in the |
| ./literal/... and ./var/... subdirectories |
| |
| See: |
| * tools/src/cmd/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 -}} |
| {{ if not (HasPrefix .Name "_") }} |
| {{- 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 enable directives */ -}} |
| {{- template "EnableDirectives" $overload -}} |
| |
| {{- /* Generate RW storage buffer parameters */ -}} |
| {{- $sb_rw_fields := Eval "EmitBufferFields" "overload" $overload |
| "var_name" "sb_rw" |
| "address_space" "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" |
| "address_space" "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" |
| "address_space" "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 "AddressSpace" $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 "AddressSpace" $p.Type -}} |
| {{- $is_abstract := DeepestElementType $p.Type | IsAbstract -}} |
| {{- if eq "function" $class -}} |
| {{- if eq "ptr" $p.Type.Target.Name -}} |
| {{template "Type" index $p.Type.TemplateArguments 1}}; |
| {{- /*indent*/}} var arg_{{$i}}: {{template "Type" index $p.Type.TemplateArguments 1}}; |
| {{ $args.Put $i (printf "&arg_%v" $i) -}} |
| {{- else if eq "var" $mode -}} |
| {{- if or $is_abstract $p.IsConst }} const arg_{{$i}} = {{Eval "ArgumentValue" $p}}; |
| {{ else }} var arg_{{$i}} = {{Eval "ArgumentValue" $p}}; |
| {{ end }} |
| {{- $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 -}} |
| ); |
| {{- /* Ensure the call isn't dead-code eliminated */ -}} |
| {{- if and $overload.ReturnType (IsDeclarable $overload.ReturnType)}} |
| {{- if IsHostShareable $overload.ReturnType}} |
| prevent_dce = res; |
| {{- else}} |
| prevent_dce = select(0, 1, all(res == {{template "Type" $overload.ReturnType}}())); |
| {{- end}} |
| {{- end}} |
| } |
| |
| {{- if and $overload.ReturnType (IsDeclarable $overload.ReturnType)}} |
| {{- if IsHostShareable $overload.ReturnType}} |
| @group(2) @binding(0) var<storage, read_write> prevent_dce : {{template "Type" $overload.ReturnType}}; |
| {{- else}} |
| @group(2) @binding(0) var<storage, read_write> prevent_dce : i32; |
| {{ end}} |
| {{ end}} |
| |
| {{- 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 "EnableDirectives" -}} |
| {{- /* Emits the required enable directives for a given overload */ -}} |
| {{- /* ------------------------------------------------------------------ */ -}} |
| {{- $permutation := . -}} |
| {{- $overload := $permutation.Overload -}} |
| {{- $builtin_name := $permutation.Intrinsic.Name -}} |
| |
| {{- /* Check and emit chromium_experimental_dp4a */ -}} |
| {{- if or (eq "dot4I8Packed" $builtin_name) (eq "dot4U8Packed" $builtin_name)}} |
| enable chromium_experimental_dp4a; |
| {{ end -}} |
| |
| {{- /* Check and emit f16 */ -}} |
| {{- if OverloadUsesF16 $overload}} |
| enable f16; |
| {{ end -}} |
| |
| {{- end -}} |
| |
| |
| {{- /* ------------------------------------------------------------------ */ -}} |
| {{- define "EmitBufferFields" -}} |
| {{- /* Emits a struct with the fields that match the given address space */ -}} |
| {{- /* 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 */ -}} |
| {{- /* 'address_space' - filtered address space */ -}} |
| {{- /* 'access' - filtered access */ -}} |
| {{- /* 'args' - argument map that's populated with the fields */ -}} |
| {{- /* ------------------------------------------------------------------ */ -}} |
| {{- $overload := .Get "overload" -}} |
| {{- $var_name := .Get "var_name" -}} |
| {{- $filter_address_space := .Get "address_space" -}} |
| {{- $filter_access := .Get "access" -}} |
| {{- $args := .Get "args" -}} |
| {{- range $i, $p := $overload.Parameters }} |
| {{- $address_space := Eval "AddressSpace" $p.Type -}} |
| {{- $access := Eval "Access" $p.Type -}} |
| {{- if and (eq $filter_address_space $address_space) (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 "AddressSpace" -}} |
| {{- /* Returns the address space 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 address space */ -}} |
| {{- /* https://gpuweb.github.io/gpuweb/wgsl/#address-space */ -}} |
| {{- $address_space := Eval "AddressSpace" . -}} |
| {{- if eq $address_space "function" -}}read_write |
| {{- else if eq $address_space "private" -}}read_write |
| {{- else if eq $address_space "workgroup" -}}read_write |
| {{- else if eq $address_space "uniform" -}}read |
| {{- else if eq $address_space "storage" -}}read |
| {{- else if eq $address_space "handle" -}}read |
| {{- end -}} |
| {{- end -}} |
| {{- end -}} |
| |
| |
| {{- /* ------------------------------------------------------------------ */ -}} |
| {{- define "ArgumentValue" -}} |
| {{- /* Returns a value that can be used for the parameter argument */ -}} |
| {{- /* ------------------------------------------------------------------ */ -}} |
| {{- $ty := .Type -}} |
| {{- $value := printf "%v" .TestValue }} |
| {{- /* $float_value is $value with a '.' suffix (if it wasn't a float already) */ -}} |
| {{- $float_value := $value }} |
| {{- if not (Contains $value ".") }}{{$float_value = printf "%v." $value}}{{end}} |
| {{- /* emit the value with any necessary suffixes */ -}} |
| {{- if eq $ty.Target.Name "i32" -}}{{$value}}i |
| {{- else if eq $ty.Target.Name "u32" -}}{{$value}}u |
| {{- else if eq $ty.Target.Name "f32" -}}{{$float_value}}f |
| {{- else if eq $ty.Target.Name "f16" -}}{{$float_value}}h |
| {{- else if eq $ty.Target.Name "fa" -}}{{$float_value}} |
| {{- else if eq $ty.Target.Name "ia" -}}{{$value}} |
| {{- else if eq $ty.Target.Name "bool" -}}true |
| {{- else if HasPrefix $ty.Target.Name "vec" -}} |
| {{- $el := Eval "ArgumentValue" "Type" (ElementType .Type) "TestValue" .TestValue }} |
| {{- template "Type" $ty}}({{$el}}) |
| {{- else if HasPrefix $ty.Target.Name "mat" -}}{{template "Type" $ty}}( |
| {{- $el := Eval "ArgumentValue" "Type" (ElementType .Type) "TestValue" .TestValue }} |
| {{- range $col := Iterate (index $ty.TemplateArguments 0) }} |
| {{- range $row := Iterate (index $ty.TemplateArguments 1) }} |
| {{- if or $col $row -}}, {{end}}{{$el}} |
| {{- end -}} |
| {{- end -}}) |
| {{- else -}}{{template "Type" $ty}}() |
| {{- end -}} |
| {{- end -}} |
| |
| |
| {{- /* ------------------------------------------------------------------ */ -}} |
| {{- define "Type" -}} |
| {{- /* Emits the WGSL for the Fully Qualified Name argument */ -}} |
| {{- /* ------------------------------------------------------------------ */ -}} |
| {{- if IsType .Target -}} |
| {{- if DeepestElementType . | IsAbstract -}} |
| {{- if eq .Target.Name "vec" -}}vec{{index .TemplateArguments 0}} |
| {{- else if eq .Target.Name "vec2" -}}vec2 |
| {{- else if eq .Target.Name "vec3" -}}vec3 |
| {{- else if eq .Target.Name "vec4" -}}vec4 |
| {{- else if eq .Target.Name "mat" -}}mat{{index .TemplateArguments 0}}x{{index .TemplateArguments 1}} |
| {{- else if HasPrefix .Target.Name "mat" -}}mat{{.Target.Name}} |
| {{- else -}} |
| {{- end -}} |
| {{- /* note: intrinsics.def has different type definitions for 'vec<N: num, T>', |
| 'vec2<T>', 'vec3<T>'. Because of this we need to check whether the type |
| name is exactly 'vec' and 'mat' and if not, then look for the prefix. */ -}} |
| {{- else if eq .Target.Name "vec" -}}vec{{index .TemplateArguments 0}}<{{template "Type" (ElementType .)}}> |
| {{- else if eq .Target.Name "mat" -}}mat{{index .TemplateArguments 0}}x{{index .TemplateArguments 1}}<{{template "Type" (ElementType .)}}> |
| {{- else if HasPrefix .Target.Name "vec" -}}{{.Target.Name}}<{{template "Type" (ElementType .)}}> |
| {{- else if HasPrefix .Target.Name "mat" -}}{{.Target.Name}}<{{template "Type" (ElementType .)}}> |
| {{- 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 -}} |
