{{- /* -------------------------------------------------------------------------------- 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//.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//.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 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 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 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 arg_{{$i}}: {{$el_type}}; {{ $args.Put $i (printf "&arg_%v" $i) -}} {{- else if eq "private" $class -}} var 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 arg_{{$i}}: {{$type}}; {{ $args.Put $i (printf "arg_%v" $i) -}} {{- else if eq "private" $class -}} var 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 prevent_dce : {{template "Type" $overload.ReturnType}}; {{- else}} @group(2) @binding(0) var prevent_dce : i32; {{ end}} {{ end}} {{- if $overload.CanBeUsedInStage.Vertex }} @vertex fn vertex_main() -> @builtin(position) vec4 { {{$permutation}}(); return vec4(); } {{ 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', 'vec2', 'vec3'. 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 -}} {{- 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 -}}