dawn-cmake/test/intrinsics/intrinsics.wgsl.tmpl

{{- /*
--------------------------------------------------------------------------------
Template file for use with tools/intrinsic-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 function... */ -}}
{{- range .Sem.Functions -}}
{{-   range .Overloads  -}}
{{-     range Permute . -}}
{{- /*    Generate a ./gen/<function>/<permuataion-hash>.wgsl file using
          the Permutation macro defined below                             */ -}}
{{-       $file := printf "./gen/%v/%v.wgsl" .Function.Name .Hash -}}
{{-       $content := Eval "Permutation" . -}}
{{-       WriteFile $file $content -}}
{{-     end  }}
{{-   end  }}
{{- end  }}


{{- /* ------------------------------------------------------------------ */ -}}
{{-                          define "Permutation"                            -}}
{{- /* Emits the body of the intrinsic permuation .wgsl file              */ -}}
{{- /* ------------------------------------------------------------------ */ -}}
{{-   $function    := .Function.Name -}}
{{-   $permutation := printf "%v_%v" $function .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 -}}
[[block]]
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 -}}
[[block]]
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 -}}
[[block]]
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: {{template "Type" .ReturnType}} = {{/* preserve space after = */ -}}
{{-   end -}}
  {{$function}}(
{{-   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 -}}