1 // Copyright 2021 The Tint Authors. 2 // 3 // Licensed under the Apache License, Version 2.0 (the "License"); 4 // you may not use this file except in compliance with the License. 5 // You may obtain a copy of the License at 6 // 7 // http://www.apache.org/licenses/LICENSE-2.0 8 // 9 // Unless required by applicable law or agreed to in writing, software 10 // distributed under the License is distributed on an "AS IS" BASIS, 11 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. 12 // See the License for the specific language governing permissions and 13 // limitations under the License. 14 15 #ifndef SRC_TRANSFORM_MODULE_SCOPE_VAR_TO_ENTRY_POINT_PARAM_H_ 16 #define SRC_TRANSFORM_MODULE_SCOPE_VAR_TO_ENTRY_POINT_PARAM_H_ 17 18 #include "src/transform/transform.h" 19 20 namespace tint { 21 namespace transform { 22 23 /// Move module-scope variables into the entry point as parameters. 24 /// 25 /// MSL does not allow module-scope variables to have any address space other 26 /// than `constant`. This transform moves all module-scope declarations into the 27 /// entry point function (either as parameters or function-scope variables) and 28 /// then passes them as pointer parameters to any function that references them. 29 /// 30 /// Since WGSL does not allow entry point parameters or function-scope variables 31 /// to have these storage classes, we annotate the new variable declarations 32 /// with an attribute that bypasses that validation rule. 33 /// 34 /// Before: 35 /// ``` 36 /// [[block]] 37 /// struct S { 38 /// f : f32; 39 /// }; 40 /// [[binding(0), group(0)]] 41 /// var<storage, read> s : S; 42 /// var<private> p : f32 = 2.0; 43 /// 44 /// fn foo() { 45 /// p = p + f; 46 /// } 47 /// 48 /// [[stage(compute), workgroup_size(1)]] 49 /// fn main() { 50 /// foo(); 51 /// } 52 /// ``` 53 /// 54 /// After: 55 /// ``` 56 /// fn foo(p : ptr<private, f32>, sptr : ptr<storage, S, read>) { 57 /// *p = *p + (*sptr).f; 58 /// } 59 /// 60 /// [[stage(compute), workgroup_size(1)]] 61 /// fn main(sptr : ptr<storage, S, read>) { 62 /// var<private> p : f32 = 2.0; 63 /// foo(&p, sptr); 64 /// } 65 /// ``` 66 class ModuleScopeVarToEntryPointParam 67 : public Castable<ModuleScopeVarToEntryPointParam, Transform> { 68 public: 69 /// Constructor 70 ModuleScopeVarToEntryPointParam(); 71 /// Destructor 72 ~ModuleScopeVarToEntryPointParam() override; 73 74 protected: 75 /// Runs the transform using the CloneContext built for transforming a 76 /// program. Run() is responsible for calling Clone() on the CloneContext. 77 /// @param ctx the CloneContext primed with the input program and 78 /// ProgramBuilder 79 /// @param inputs optional extra transform-specific input data 80 /// @param outputs optional extra transform-specific output data 81 void Run(CloneContext& ctx, const DataMap& inputs, DataMap& outputs) override; 82 83 struct State; 84 }; 85 86 } // namespace transform 87 } // namespace tint 88 89 #endif // SRC_TRANSFORM_MODULE_SCOPE_VAR_TO_ENTRY_POINT_PARAM_H_ 90