{{- /* -------------------------------------------------------------------------------- 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 -}} ); } {{/*new line*/ -}} {{- 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 "fa" -}}{{$float_value}} {{- else if eq $ty.Target.Name "ia" -}}{{$value}} {{- else if eq $ty.Target.Name "bool" -}}true {{- else if eq $ty.Target.Name "vec" -}} {{- $el := Eval "ArgumentValue" "Type" (ElementType .Type) "Name" "" "TestValue" .TestValue }} {{- template "Type" $ty}}({{$el}}) {{- else if eq $ty.Target.Name "mat" -}}{{template "Type" $ty}}( {{- $el := Eval "ArgumentValue" "Type" (ElementType .Type) "Name" "" "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 "mat" -}}mat{{index .TemplateArguments 0}}x{{index .TemplateArguments 1}} {{- else -}} {{- end -}} {{- 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 -}}{{.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 -}}