1{{- /* 2-------------------------------------------------------------------------------- 3Template file for use with tools/intrinsic-gen to generate the wgsl files in the 4./gen/... subdirectories 5 6See: 7* tools/cmd/intrinsic-gen/gen for structures used by this template 8* https://golang.org/pkg/text/template/ for documentation on the template syntax 9-------------------------------------------------------------------------------- 10*/ -}} 11 12{{- /* For each permutation of each overload of each function... */ -}} 13{{- range .Sem.Functions -}} 14{{- range .Overloads -}} 15{{- range Permute . -}} 16{{- /* Generate a ./gen/<function>/<permuataion-hash>.wgsl file using 17 the Permutation macro defined below */ -}} 18{{- $file := printf "./gen/%v/%v.wgsl" .Function.Name .Hash -}} 19{{- $content := Eval "Permutation" . -}} 20{{- WriteFile $file $content -}} 21{{- end }} 22{{- end }} 23{{- end }} 24 25 26{{- /* ------------------------------------------------------------------ */ -}} 27{{- define "Permutation" -}} 28{{- /* Emits the body of the intrinsic permuation .wgsl file */ -}} 29{{- /* ------------------------------------------------------------------ */ -}} 30{{- $function := .Function.Name -}} 31{{- $permutation := printf "%v_%v" $function .Hash -}} 32{{- $args := Map -}} 33 34{{- /* Generate RW storage buffer parameters */ -}} 35{{- $sb_rw_fields := Eval "EmitBufferFields" "overload" . 36 "var_name" "sb_rw" 37 "storage" "storage" 38 "access" "read_write" 39 "args" $args -}} 40{{- if $sb_rw_fields -}} 41[[block]] 42struct SB_RW { 43{{- $sb_rw_fields -}} 44}; 45[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW; 46{{ end -}} 47 48{{- /* Generate RO storage buffer parameters */ -}} 49{{- $sb_ro_fields := Eval "EmitBufferFields" "overload" . 50 "var_name" "sb_ro" 51 "storage" "storage" 52 "access" "read" 53 "args" $args -}} 54{{- if $sb_ro_fields -}} 55[[block]] 56struct SB_RO { 57{{- $sb_ro_fields -}} 58}; 59[[group(0), binding(1)]] var<storage, read> sb_ro : SB_RO; 60{{ end -}} 61 62{{- /* Generate uniform buffer parameters */ -}} 63{{- $ub_fields := Eval "EmitBufferFields" "overload" . 64 "var_name" "ub" 65 "storage" "uniform" 66 "access" "read" 67 "args" $args -}} 68{{- if $ub_fields -}} 69[[block]] 70struct UB { 71{{- $ub_fields -}} 72}; 73[[group(0), binding(1)]] var<uniform> ub : UB; 74{{ end -}} 75 76{{- /* Generate module-scoped handle variables */ -}} 77{{- range $i, $p := .Parameters }} 78{{- $class := Eval "StorageClass" $p.Type -}} 79{{- if eq "ptr" $p.Type.Target.Name -}} 80{{- $el_type := Eval "Type" (index $p.Type.TemplateArguments 1)}} 81{{- if eq "handle" $class -}} 82 [[group(1), binding({{$i}})]] var arg_{{$i}}: {{$el_type}}; 83{{ $args.Put $i (printf "&arg_%v" $i) -}} 84{{- else if eq "workgroup" $class -}} 85 var<workgroup> arg_{{$i}}: {{$el_type}}; 86{{ $args.Put $i (printf "&arg_%v" $i) -}} 87{{- else if eq "private" $class -}} 88 var<private> arg_{{$i}}: {{$el_type}}; 89{{ $args.Put $i (printf "&arg_%v" $i) -}} 90{{- end -}} 91{{- else -}} 92{{- $type := Eval "Type" $p.Type}} 93{{- if eq "handle" $class -}} 94 [[group(1), binding({{$i}})]] var arg_{{$i}}: {{$type}}; 95{{ $args.Put $i (printf "arg_%v" $i) -}} 96{{- else if eq "workgroup" $class -}} 97 var<workgroup> arg_{{$i}}: {{$type}}; 98{{ $args.Put $i (printf "arg_%v" $i) -}} 99{{- else if eq "private" $class -}} 100 var<private> arg_{{$i}}: {{$type}}; 101{{ $args.Put $i (printf "arg_%v" $i) -}} 102{{- end -}} 103{{- end -}} 104{{- end -}} 105 106{{- /* Generate the function that calls the intrinsic */ -}} 107{{- /*newline*/}} 108// {{$.Overload}} 109fn {{$permutation}}() { 110{{/* Build the parameters either as 'var' or inline values */ -}} 111{{- range $i, $p := .Parameters -}} 112{{- $class := Eval "StorageClass" $p.Type -}} 113{{- if eq "function" $class -}} 114{{- if eq "ptr" $p.Type.Target.Name -}} 115{{- /*indent*/}} var arg_{{$i}}: {{template "Type" index $p.Type.TemplateArguments 1}}; 116{{ $args.Put $i (printf "&arg_%v" $i) -}} 117{{- else -}} 118{{- $args.Put $i (Eval "ArgumentValue" $p) -}} 119{{- end -}} 120{{- end -}} 121{{- end -}} 122 123{{- /* Make the call to the intrinsic */ -}} 124{{- /*indent*/}} {{/*indent*/ -}} 125{{- if .ReturnType -}} 126 var res{{if IsDeclarable .ReturnType}}: {{template "Type" .ReturnType}}{{end}} = {{/* preserve space after = */ -}} 127{{- end -}} 128 {{$function}}( 129{{- range $i, $p := .Parameters -}} 130{{- if $i -}}, {{end}}{{$args.Get $i -}} 131{{- end -}} 132 ); 133} 134{{/*new line*/ -}} 135 136{{- if .CanBeUsedInStage.Vertex }} 137[[stage(vertex)]] 138fn vertex_main() -> [[builtin(position)]] vec4<f32> { 139 {{$permutation}}(); 140 return vec4<f32>(); 141} 142{{ end -}} 143 144{{- if .CanBeUsedInStage.Fragment }} 145[[stage(fragment)]] 146fn fragment_main() { 147 {{$permutation}}(); 148} 149{{ end -}} 150 151{{- if .CanBeUsedInStage.Compute }} 152[[stage(compute), workgroup_size(1)]] 153fn compute_main() { 154 {{$permutation}}(); 155} 156{{ end -}} 157 158{{- end -}} 159 160 161{{- /* ------------------------------------------------------------------ */ -}} 162{{- define "EmitBufferFields" -}} 163{{- /* Emits a struct with the fields that match the given storage class */ -}} 164{{- /* and access. */ -}} 165{{- /* Argument is a map with the following expected keys: */ -}} 166{{- /* 'overload' - the current overload */ -}} 167{{- /* 'var_name' - name of the variable of the structure type */ -}} 168{{- /* 'storage' - filtered storage class */ -}} 169{{- /* 'access' - filtered access */ -}} 170{{- /* 'args' - argument map that's populated with the fields */ -}} 171{{- /* ------------------------------------------------------------------ */ -}} 172{{- $overload := .Get "overload" -}} 173{{- $var_name := .Get "var_name" -}} 174{{- $filter_storage := .Get "storage" -}} 175{{- $filter_access := .Get "access" -}} 176{{- $args := .Get "args" -}} 177{{- range $i, $p := $overload.Parameters }} 178{{- $storage := Eval "StorageClass" $p.Type -}} 179{{- $access := Eval "Access" $p.Type -}} 180{{- if and (eq $filter_storage $storage) (eq $filter_access $access) }} 181{{- if eq "ptr" $p.Type.Target.Name }} 182 arg_{{$i}}: {{template "Type" (index $p.Type.TemplateArguments 1)}}; 183{{ $args.Put $i (printf "&%v.arg_%v" $var_name $i) -}} 184{{- else }} 185 arg_{{$i}}: {{template "Type" $p.Type}}; 186{{ $args.Put $i (printf "%v.arg_%v" $var_name $i) -}} 187{{- end -}} 188{{- end -}} 189{{- end -}} 190{{ end -}} 191 192{{- /* ------------------------------------------------------------------ */ -}} 193{{- define "StorageClass" -}} 194{{- /* Returns the storage class for the given Fully Qualified Name */ -}} 195{{- /* ------------------------------------------------------------------ */ -}} 196{{- $name := .Target.Name -}} 197{{- if eq $name "array" -}}storage 198{{- else if HasPrefix $name "texture" -}}handle 199{{- else if HasPrefix $name "sampler" -}}handle 200{{- else if eq $name "ptr" -}}{{(index .TemplateArguments 0).Target.Name}} 201{{- else -}}function 202{{- end -}} 203{{- end -}} 204 205 206{{- /* ------------------------------------------------------------------ */ -}} 207{{- define "Access" -}} 208{{- /* Returns the access for the given Fully Qualified Name */ -}} 209{{- /* ------------------------------------------------------------------ */ -}} 210{{- $name := .Target.Name -}} 211{{- if eq $name "ptr" -}}{{(index .TemplateArguments 2).Target.Name}} 212{{- else -}} 213{{- /* Emit the default for the storage class */ -}} 214{{- /* https://gpuweb.github.io/gpuweb/wgsl/#storage-class */ -}} 215{{- $storage := Eval "StorageClass" . -}} 216{{- if eq $storage "function" -}}read_write 217{{- else if eq $storage "private" -}}read_write 218{{- else if eq $storage "workgroup" -}}read_write 219{{- else if eq $storage "uniform" -}}read 220{{- else if eq $storage "storage" -}}read 221{{- else if eq $storage "handle" -}}read 222{{- end -}} 223{{- end -}} 224{{- end -}} 225 226 227{{- /* ------------------------------------------------------------------ */ -}} 228{{- define "ArgumentValue" -}} 229{{- /* Returns a value that can be used for the parameter argument */ -}} 230{{- /* ------------------------------------------------------------------ */ -}} 231{{- $ty := .Type -}} 232{{- if eq $ty.Target.Name "i32" -}} 233{{- /* If the parameter has the name 'level', then use '0' as the value. */ -}} 234{{- /* Some texture arguments require the level parameter to be 0, and */ -}} 235{{- /* constraint is not described in the definition file. */ -}} 236{{- if eq .Name "level" -}}0 237{{- else -}}1 238{{- end -}} 239{{- else if eq $ty.Target.Name "u32" -}}1u 240{{- else if eq $ty.Target.Name "f32" -}}1.0 241{{- else -}}{{template "Type" $ty}}() 242{{- end -}} 243{{- end -}} 244 245 246{{- /* ------------------------------------------------------------------ */ -}} 247{{- define "Type" -}} 248{{- /* Emits the WGSL for the Fully Qualified Name argument */ -}} 249{{- /* ------------------------------------------------------------------ */ -}} 250{{- if IsType .Target -}} 251{{- if eq .Target.Name "vec" -}}vec{{index .TemplateArguments 0}}<{{template "Type" index .TemplateArguments 1}}> 252{{- else if eq .Target.Name "mat" -}}mat{{index .TemplateArguments 0}}x{{index .TemplateArguments 1}}<{{template "Type" index .TemplateArguments 2}}> 253{{- else -}}{{.Target.Name}}{{template "TemplateArguments" .TemplateArguments}} 254{{- end -}} 255{{- else if IsEnumEntry .Target -}}{{.Target.Name}} 256{{- else if IsEnumMatcher .Target -}}{{(index .Target.Options 0).Name}} 257{{- else -}}<unhandled-fully-qualified-name-target={{- printf "%T" .Target -}}> 258{{- end -}} 259{{- end -}} 260 261 262{{- /* ------------------------------------------------------------------ */ -}} 263{{- define "TemplateArguments" -}} 264{{- /* Emits the WGSL for the template argument list */ -}} 265{{- /* ------------------------------------------------------------------ */ -}} 266{{- if . -}} 267< 268{{- range $i, $a := . -}} 269{{- if $i -}}, {{ end -}} 270{{- if IsInt $a -}}{{- . -}} 271{{- else -}}{{- template "Type" $a -}} 272{{- end -}} 273{{- end -}} 274> 275{{- end -}} 276{{- end -}} 277