encounter/dawn-cmake
1
0
Fork 0
mirror of https://github.com/encounter/dawn-cmake.git synced 2025-05-13 10:51:35 +00:00
Code Issues Packages Projects Releases Wiki Activity

Update builtin methods.

Browse Source 
This CL removes the import mechanism, the identifier paths and updates
all of the standard library methods to be builtins.

Bug: tint:242
Change-Id: If09b98a155ae49ced3986ba2c9b517a060693006
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/28720
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: Ryan Harrison <rharrison@chromium.org>
This commit is contained in:
dan sinclair 2020-09-22 19:42:13 +00:00 committed by Commit Bot service account
parent 5f8126271d
commit b4fee2f824
53 changed files with 3959 additions and 2866 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
BUILD.gn
View File

@ -269,8 +269,6 @@ source_set("libtint_core_src") {
"src/ast/identifier_expression.h",
"src/ast/if_statement.cc",
"src/ast/if_statement.h",
"src/ast/import.cc",
"src/ast/import.h",
"src/ast/int_literal.cc",
"src/ast/int_literal.h",
"src/ast/intrinsic.cc",
@ -710,7 +708,6 @@ source_set("tint_unittests_core_src") {
"src/ast/function_test.cc",
"src/ast/identifier_expression_test.cc",
"src/ast/if_statement_test.cc",
"src/ast/import_test.cc",
"src/ast/int_literal_test.cc",
"src/ast/location_decoration_test.cc",
"src/ast/loop_statement_test.cc",
@ -908,7 +905,6 @@ source_set("tint_unittests_wgsl_reader_src") {
"src/reader/wgsl/parser_impl_global_variable_decl_test.cc",
"src/reader/wgsl/parser_impl_if_stmt_test.cc",
"src/reader/wgsl/parser_impl_image_storage_type_test.cc",
"src/reader/wgsl/parser_impl_import_decl_test.cc",
"src/reader/wgsl/parser_impl_inclusive_or_expression_test.cc",
"src/reader/wgsl/parser_impl_logical_and_expression_test.cc",
"src/reader/wgsl/parser_impl_logical_or_expression_test.cc",
@ -991,7 +987,6 @@ source_set("tint_unittests_wgsl_writer_src") {
"src/writer/wgsl/generator_impl_function_test.cc",
"src/writer/wgsl/generator_impl_identifier_test.cc",
"src/writer/wgsl/generator_impl_if_test.cc",
"src/writer/wgsl/generator_impl_import_test.cc",
"src/writer/wgsl/generator_impl_loop_test.cc",
"src/writer/wgsl/generator_impl_member_accessor_test.cc",
"src/writer/wgsl/generator_impl_return_test.cc",

5
src/CMakeLists.txt
View File

@ -90,8 +90,6 @@ set(TINT_LIB_SRCS
ast/identifier_expression.h
ast/if_statement.cc
ast/if_statement.h
ast/import.cc
ast/import.h
ast/int_literal.cc
ast/int_literal.h
ast/intrinsic.cc
@ -319,7 +317,6 @@ set(TINT_TEST_SRCS
ast/function_test.cc
ast/identifier_expression_test.cc
ast/if_statement_test.cc
ast/import_test.cc
ast/int_literal_test.cc
ast/location_decoration_test.cc
ast/loop_statement_test.cc
@ -439,7 +436,6 @@ if(${TINT_BUILD_WGSL_READER})
reader/wgsl/parser_impl_global_variable_decl_test.cc
reader/wgsl/parser_impl_if_stmt_test.cc
reader/wgsl/parser_impl_image_storage_type_test.cc
reader/wgsl/parser_impl_import_decl_test.cc
reader/wgsl/parser_impl_inclusive_or_expression_test.cc
reader/wgsl/parser_impl_logical_and_expression_test.cc
reader/wgsl/parser_impl_logical_or_expression_test.cc
@ -538,7 +534,6 @@ if(${TINT_BUILD_WGSL_WRITER})
writer/wgsl/generator_impl_function_test.cc
writer/wgsl/generator_impl_identifier_test.cc
writer/wgsl/generator_impl_if_test.cc
writer/wgsl/generator_impl_import_test.cc
writer/wgsl/generator_impl_loop_test.cc
writer/wgsl/generator_impl_member_accessor_test.cc
writer/wgsl/generator_impl_return_test.cc

43
src/ast/identifier_expression.cc
View File

@ -18,62 +18,27 @@ namespace tint {
namespace ast {
IdentifierExpression::IdentifierExpression(const std::string& name)
: Expression(), segments_({name}) {}
: Expression(), name_(name) {}
IdentifierExpression::IdentifierExpression(const Source& source,
const std::string& name)
: Expression(source), segments_({name}) {}
IdentifierExpression::IdentifierExpression(std::vector<std::string> segments)
: Expression(), segments_(std::move(segments)) {}
IdentifierExpression::IdentifierExpression(const Source& source,
std::vector<std::string> segments)
: Expression(source), segments_(std::move(segments)) {}
: Expression(source), name_(name) {}
IdentifierExpression::IdentifierExpression(IdentifierExpression&&) = default;
IdentifierExpression::~IdentifierExpression() = default;
std::string IdentifierExpression::path() const {
if (segments_.size() < 2) {
return "";
}
std::string path = "";
// We skip the last segment as that's the name, not part of the path
for (size_t i = 0; i < segments_.size() - 1; ++i) {
if (i > 0) {
path += "::";
}
path += segments_[i];
}
return path;
}
bool IdentifierExpression::IsIdentifier() const {
return true;
}
bool IdentifierExpression::IsValid() const {
if (segments_.size() == 0)
return false;
for (const auto& segment : segments_) {
if (segment.size() == 0)
return false;
}
return true;
return !name_.empty();
}
void IdentifierExpression::to_str(std::ostream& out, size_t indent) const {
make_indent(out, indent);
out << "Identifier{";
if (has_path()) {
out << path() << "::";
}
out << name();
out << "}" << std::endl;
out << "Identifier{" << name_ << "}" << std::endl;
}
} // namespace ast

35
src/ast/identifier_expression.h
View File

@ -16,10 +16,9 @@
#define SRC_AST_IDENTIFIER_EXPRESSION_H_
#include <string>
#include <utility>
#include <vector>
#include "src/ast/expression.h"
#include "src/ast/intrinsic.h"
namespace tint {
namespace ast {
@ -34,31 +33,20 @@ class IdentifierExpression : public Expression {
/// @param source the source
/// @param name the name
IdentifierExpression(const Source& source, const std::string& name);
/// Constructor
/// @param segments the name segments
explicit IdentifierExpression(std::vector<std::string> segments);
/// Constructor
/// @param source the identifier expression source
/// @param segments the name segments
IdentifierExpression(const Source& source, std::vector<std::string> segments);
/// Move constructor
IdentifierExpression(IdentifierExpression&&);
~IdentifierExpression() override;
/// Sets the name segments
/// @param segments the name segments
void set_segments(std::vector<std::string> segments) {
segments_ = std::move(segments);
}
/// @returns the name
std::vector<std::string> segments() const { return segments_; }
/// @returns true if this identifier has a path component
bool has_path() const { return segments_.size() > 1; }
/// @returns the path part of the identifier or blank if no path
std::string path() const;
/// @returns the name part of the identifier
std::string name() const { return segments_.back(); }
std::string name() const { return name_; }
/// Sets the intrinsic for this identifier
/// @param i the intrinsic to set
void set_intrinsic(Intrinsic i) { intrinsic_ = i; }
/// @returns the intrinsic this identifier represents
Intrinsic intrinsic() const { return intrinsic_; }
/// @returns true if this identifier is for an intrinsic
bool IsIntrinsic() const { return intrinsic_ != Intrinsic::kNone; }
/// @returns true if this is an identifier expression
bool IsIdentifier() const override;
@ -74,7 +62,8 @@ class IdentifierExpression : public Expression {
private:
IdentifierExpression(const IdentifierExpression&) = delete;
std::vector<std::string> segments_;
Intrinsic intrinsic_ = Intrinsic::kNone;
std::string name_;
};
} // namespace ast

28
src/ast/identifier_expression_test.cc
View File

@ -24,19 +24,11 @@ using IdentifierExpressionTest = testing::Test;
TEST_F(IdentifierExpressionTest, Creation) {
IdentifierExpression i("ident");
ASSERT_EQ(i.segments().size(), 1u);
EXPECT_EQ(i.segments()[0], "ident");
EXPECT_EQ(i.path(), "");
EXPECT_EQ(i.name(), "ident");
}
TEST_F(IdentifierExpressionTest, Creation_WithSource) {
IdentifierExpression i(Source{20, 2}, {"ns1", "ns2", "ident"});
ASSERT_EQ(i.segments().size(), 3u);
EXPECT_EQ(i.segments()[0], "ns1");
EXPECT_EQ(i.segments()[1], "ns2");
EXPECT_EQ(i.segments()[2], "ident");
EXPECT_EQ(i.path(), "ns1::ns2");
IdentifierExpression i(Source{20, 2}, "ident");
EXPECT_EQ(i.name(), "ident");
auto src = i.source();
@ -54,21 +46,11 @@ TEST_F(IdentifierExpressionTest, IsValid) {
EXPECT_TRUE(i.IsValid());
}
TEST_F(IdentifierExpressionTest, IsValid_WithNamespaces) {
IdentifierExpression i({"ns1", "n2", "ident"});
EXPECT_TRUE(i.IsValid());
}
TEST_F(IdentifierExpressionTest, IsValid_BlankName) {
IdentifierExpression i("");
EXPECT_FALSE(i.IsValid());
}
TEST_F(IdentifierExpressionTest, IsValid_BlankNamespace) {
IdentifierExpression i({"ns1", "", "ident"});
EXPECT_FALSE(i.IsValid());
}
TEST_F(IdentifierExpressionTest, ToStr) {
IdentifierExpression i("ident");
std::ostringstream out;
@ -77,14 +59,6 @@ TEST_F(IdentifierExpressionTest, ToStr) {
)");
}
TEST_F(IdentifierExpressionTest, ToStr_WithNamespace) {
IdentifierExpression i({"ns1", "ns2", "ident"});
std::ostringstream out;
i.to_str(out, 2);
EXPECT_EQ(out.str(), R"( Identifier{ns1::ns2::ident}
)");
}
} // namespace
} // namespace ast
} // namespace tint

60
src/ast/import.cc
View File

@ -1,60 +0,0 @@
// Copyright 2020 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include "src/ast/import.h"
#include <cctype>
namespace tint {
namespace ast {
Import::Import() = default;
Import::Import(const std::string& path, const std::string& name)
: Node(), path_(path), name_(name) {}
Import::Import(const Source& source,
const std::string& path,
const std::string& name)
: Node(source), path_(path), name_(name) {}
Import::Import(Import&&) = default;
Import::~Import() = default;
bool Import::IsValid() const {
if (path_.length() == 0) {
return false;
}
auto len = name_.length();
if (len == 0) {
return false;
}
// Verify the import name ends in a character, number or _
if (len > 2 && !std::isalnum(name_[len - 1]) && name_[len] != '_') {
return false;
}
return true;
}
void Import::to_str(std::ostream& out, size_t indent) const {
make_indent(out, indent);
out << R"(Import{")" + path_ + R"(" as )" + name_ + "}" << std::endl;
}
} // namespace ast
} // namespace tint

101
src/ast/import.h
View File

@ -1,101 +0,0 @@
// Copyright 2020 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#ifndef SRC_AST_IMPORT_H_
#define SRC_AST_IMPORT_H_
#include <memory>
#include <ostream>
#include <string>
#include <unordered_map>
#include <vector>
#include "src/ast/node.h"
namespace tint {
namespace ast {
/// An import statement.
class Import : public Node {
public:
/// Create a new empty import statement
Import();
/// Create a new import statement
/// @param path The import path e.g. GLSL.std.450
/// @param name The import reference name e.g. std
Import(const std::string& path, const std::string& name);
/// Create a new import statement
/// @param source The input source for the import statement
/// @param path The import path e.g. GLSL.std.430
/// @param name The import reference name e.g. std
Import(const Source& source,
const std::string& path,
const std::string& name);
/// Move constructor
Import(Import&&);
~Import() override;
/// Sets the import path
/// @param path the path to set
void set_path(const std::string& path) { path_ = path; }
/// @returns the import path
const std::string& path() const { return path_; }
/// Sets the import name
/// @param name the name to set
void set_name(const std::string& name) { name_ = name; }
/// @returns the import name
const std::string& name() const { return name_; }
/// Add the given |name| to map to the given |id|
/// @param name the name to map
/// @param id the id to map too
void AddMethodId(const std::string& name, uint32_t id) {
method_to_id_map_[name] = id;
}
/// Retrieves the id for a given name
/// @param name the name to lookup
/// @returns the id for the given name or 0 on failure
uint32_t GetIdForMethod(const std::string& name) const {
auto val = method_to_id_map_.find(name);
if (val == method_to_id_map_.end()) {
return 0;
}
return val->second;
}
/// @returns true if the name and path are both present
bool IsValid() const override;
/// Writes a representation of the node to the output stream
/// @param out the stream to write to
/// @param indent number of spaces to indent the node when writing
void to_str(std::ostream& out, size_t indent) const override;
private:
Import(const Import&) = delete;
std::string path_;
std::string name_;
std::unordered_map<std::string, uint32_t> method_to_id_map_;
};
/// A list of unique imports
using ImportList = std::vector<std::unique_ptr<Import>>;
} // namespace ast
} // namespace tint
#endif // SRC_AST_IMPORT_H_

93
src/ast/import_test.cc
View File

@ -1,93 +0,0 @@
// Copyright 2020 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include "src/ast/import.h"
#include <sstream>
#include "gtest/gtest.h"
namespace tint {
namespace ast {
namespace {
using ImportTest = testing::Test;
TEST_F(ImportTest, Creation) {
Import i("GLSL.std.430", "std::glsl");
EXPECT_EQ(i.path(), "GLSL.std.430");
EXPECT_EQ(i.name(), "std::glsl");
EXPECT_EQ(i.line(), 0u);
EXPECT_EQ(i.column(), 0u);
}
TEST_F(ImportTest, CreationWithSource) {
Source s{27, 4};
Import i(s, "GLSL.std.430", "std::glsl");
EXPECT_EQ(i.path(), "GLSL.std.430");
EXPECT_EQ(i.name(), "std::glsl");
EXPECT_EQ(i.line(), 27u);
EXPECT_EQ(i.column(), 4u);
}
TEST_F(ImportTest, CreationEmpty) {
Source s{27, 4};
Import i;
i.set_source(s);
i.set_path("GLSL.std.430");
i.set_name("std::glsl");
EXPECT_EQ(i.path(), "GLSL.std.430");
EXPECT_EQ(i.name(), "std::glsl");
EXPECT_EQ(i.line(), 27u);
EXPECT_EQ(i.column(), 4u);
}
TEST_F(ImportTest, to_str) {
Import i{"GLSL.std.430", "std::glsl"};
std::ostringstream out;
i.to_str(out, 2);
EXPECT_EQ(out.str(), " Import{\"GLSL.std.430\" as std::glsl}\n");
}
TEST_F(ImportTest, IsValid) {
Import i{"GLSL.std.430", "std::glsl"};
EXPECT_TRUE(i.IsValid());
}
TEST_F(ImportTest, IsValid_MissingPath) {
Import i{"", "std::glsl"};
EXPECT_FALSE(i.IsValid());
}
TEST_F(ImportTest, IsValid_MissingName) {
Import i{"GLSL.std.430", ""};
EXPECT_FALSE(i.IsValid());
}
TEST_F(ImportTest, IsValid_MissingBoth) {
Import i;
EXPECT_FALSE(i.IsValid());
}
TEST_F(ImportTest, IsValid_InvalidEndingCharacter) {
Import i{"GLSL.std.430", "std::glsl::"};
EXPECT_FALSE(i.IsValid());
}
} // namespace
} // namespace ast
} // namespace tint

247
src/ast/intrinsic.cc
View File

@ -16,37 +16,242 @@
namespace tint {
namespace ast {
std::ostream& operator<<(std::ostream& out, Intrinsic i) {
switch (i) {
case Intrinsic::kAbs:
out << "abs";
break;
case Intrinsic::kAcos:
out << "acos";
break;
case Intrinsic::kAll:
out << "all";
break;
case Intrinsic::kAny:
out << "any";
break;
case Intrinsic::kAsin:
out << "asin";
break;
case Intrinsic::kAtan:
out << "atan";
break;
case Intrinsic::kAtan2:
out << "atan2";
break;
case Intrinsic::kCeil:
out << "ceil";
break;
case Intrinsic::kClamp:
out << "clamp";
break;
case Intrinsic::kCos:
out << "cos";
break;
case Intrinsic::kCosh:
out << "cosh";
break;
case Intrinsic::kCountOneBits:
out << "countOneBits";
break;
case Intrinsic::kCross:
out << "cross";
break;
case Intrinsic::kDeterminant:
out << "determinant";
break;
case Intrinsic::kDistance:
out << "distance";
break;
case Intrinsic::kDot:
out << "dot";
break;
case Intrinsic::kDpdx:
out << "dpdx";
break;
case Intrinsic::kDpdxCoarse:
out << "dpdxCoarse";
break;
case Intrinsic::kDpdxFine:
out << "dpdxFine";
break;
case Intrinsic::kDpdy:
out << "dpdy";
break;
case Intrinsic::kDpdyCoarse:
out << "dpdyCoarse";
break;
case Intrinsic::kDpdyFine:
out << "dpdyFine";
break;
case Intrinsic::kExp:
out << "exp";
break;
case Intrinsic::kExp2:
out << "exp2";
break;
case Intrinsic::kFaceForward:
out << "faceForward";
break;
case Intrinsic::kFloor:
out << "floor";
break;
case Intrinsic::kFma:
out << "fma";
break;
case Intrinsic::kFract:
out << "fract";
break;
case Intrinsic::kFrexp:
out << "frexp";
break;
case Intrinsic::kFwidth:
out << "fwidth";
break;
case Intrinsic::kFwidthCoarse:
out << "fwidthCoarse";
break;
case Intrinsic::kFwidthFine:
out << "fwidthFine";
break;
case Intrinsic::kInverseSqrt:
out << "inverseSqrt";
break;
case Intrinsic::kIsFinite:
out << "isFinite";
break;
case Intrinsic::kIsInf:
out << "isInf";
break;
case Intrinsic::kIsNan:
out << "isNan";
break;
case Intrinsic::kIsNormal:
out << "isNormal";
break;
case Intrinsic::kLdexp:
out << "ldexp";
break;
case Intrinsic::kLength:
out << "length";
break;
case Intrinsic::kLog:
out << "log";
break;
case Intrinsic::kLog2:
out << "log2";
break;
case Intrinsic::kMax:
out << "max";
break;
case Intrinsic::kMin:
out << "min";
break;
case Intrinsic::kMix:
out << "mix";
break;
case Intrinsic::kModf:
out << "modf";
break;
case Intrinsic::kNormalize:
out << "normalize";
break;
case Intrinsic::kOuterProduct:
out << "outerProduct";
break;
case Intrinsic::kPow:
out << "pow";
break;
case Intrinsic::kReflect:
out << "reflect";
break;
case Intrinsic::kReverseBits:
out << "reverseBits";
break;
case Intrinsic::kRound:
out << "round";
break;
case Intrinsic::kSelect:
out << "select";
break;
case Intrinsic::kSign:
out << "sign";
break;
case Intrinsic::kSin:
out << "sin";
break;
case Intrinsic::kSinh:
out << "sinh";
break;
case Intrinsic::kSmoothStep:
out << "smoothStep";
break;
case Intrinsic::kSqrt:
out << "sqrt";
break;
case Intrinsic::kStep:
out << "step";
break;
case Intrinsic::kTan:
out << "tan";
break;
case Intrinsic::kTanh:
out << "tanh";
break;
case Intrinsic::kTextureLoad:
out << "textureLoad";
break;
case Intrinsic::kTextureSample:
out << "textureSample";
break;
case Intrinsic::kTextureSampleBias:
out << "textureSampleBias";
break;
case Intrinsic::kTextureSampleCompare:
out << "textureSampleCompare";
break;
case Intrinsic::kTextureSampleLevel:
out << "textureSampleLevel";
break;
case Intrinsic::kTrunc:
out << "trunc";
break;
default:
out << "Unknown";
break;
}
return out;
}
namespace intrinsic {
bool IsCoarseDerivative(const std::string& name) {
return name == "dpdxCoarse" || name == "dpdyCoarse" || name == "fwidthCoarse";
bool IsCoarseDerivative(ast::Intrinsic i) {
return i == Intrinsic::kDpdxCoarse ||
i == Intrinsic::kDpdyCoarse | i == Intrinsic::kFwidthCoarse;
}
bool IsFineDerivative(const std::string& name) {
return name == "dpdxFine" || name == "dpdyFine" || name == "fwidthFine";
bool IsFineDerivative(ast::Intrinsic i) {
return i == Intrinsic::kDpdxFine || i == Intrinsic::kDpdyFine ||
i == Intrinsic::kFwidthFine;
}
bool IsDerivative(const std::string& name) {
return name == "dpdx" || name == "dpdy" || name == "fwidth" ||
IsCoarseDerivative(name) || IsFineDerivative(name);
bool IsDerivative(ast::Intrinsic i) {
return i == Intrinsic::kDpdx || i == Intrinsic::kDpdy ||
i == Intrinsic::kFwidth || IsCoarseDerivative(i) ||
IsFineDerivative(i);
}
bool IsFloatClassificationIntrinsic(const std::string& name) {
return name == "isFinite" || name == "isInf" || name == "isNan" ||
name == "isNormal";
bool IsFloatClassificationIntrinsic(ast::Intrinsic i) {
return i == Intrinsic::kIsFinite || i == Intrinsic::kIsInf ||
i == Intrinsic::kIsNan || i == Intrinsic::kIsNormal;
}
bool IsTextureOperationIntrinsic(const std::string& name) {
return name == "textureLoad" || name == "textureSample" ||
name == "textureSampleLevel" || name == "textureSampleBias" ||
name == "textureSampleCompare";
}
bool IsIntrinsic(const std::string& name) {
return IsDerivative(name) || name == "all" || name == "any" ||
IsFloatClassificationIntrinsic(name) ||
IsTextureOperationIntrinsic(name) || name == "dot" ||
name == "outerProduct" || name == "select";
bool IsTextureIntrinsic(ast::Intrinsic i) {
return i == Intrinsic::kTextureLoad || i == Intrinsic::kTextureSample ||
i == Intrinsic::kTextureSampleLevel ||
i == Intrinsic::kTextureSampleBias ||
i == Intrinsic::kTextureSampleCompare;
}
} // namespace intrinsic

100
src/ast/intrinsic.h
View File

@ -15,41 +15,111 @@
#ifndef SRC_AST_INTRINSIC_H_
#define SRC_AST_INTRINSIC_H_
#include <ostream>
#include <string>
namespace tint {
namespace ast {
enum class Intrinsic {
kNone = -1,
kAbs,
kAcos,
kAll,
kAny,
kAsin,
kAtan,
kAtan2,
kCeil,
kClamp,
kCos,
kCosh,
kCountOneBits,
kCross,
kDeterminant,
kDistance,
kDot,
kDpdx,
kDpdxCoarse,
kDpdxFine,
kDpdy,
kDpdyCoarse,
kDpdyFine,
kExp,
kExp2,
kFaceForward,
kFloor,
kFma,
kFract,
kFrexp,
kFwidth,
kFwidthCoarse,
kFwidthFine,
kInverseSqrt,
kIsFinite,
kIsInf,
kIsNan,
kIsNormal,
kLdexp,
kLength,
kLog,
kLog2,
kMax,
kMin,
kMix,
kModf,
kNormalize,
kOuterProduct,
kPow,
kReflect,
kReverseBits,
kRound,
kSelect,
kSign,
kSin,
kSinh,
kSmoothStep,
kSqrt,
kStep,
kTan,
kTanh,
kTextureLoad,
kTextureSample,
kTextureSampleBias,
kTextureSampleCompare,
kTextureSampleLevel,
kTrunc
};
std::ostream& operator<<(std::ostream& out, Intrinsic i);
namespace intrinsic {
/// Determines if the given |name| is a coarse derivative
/// @param name the name to check
/// @param i the intrinsic
/// @returns true if the given derivative is coarse.
bool IsCoarseDerivative(const std::string& name);
bool IsCoarseDerivative(ast::Intrinsic i);
/// Determines if the given |name| is a fine derivative
/// @param name the name to check
/// @param i the intrinsic
/// @returns true if the given derivative is fine.
bool IsFineDerivative(const std::string& name);
bool IsFineDerivative(ast::Intrinsic i);
/// Determine if the given |name| is a derivative intrinsic
/// @param name the name to check
/// @param i the intrinsic
/// @returns true if the given |name| is a derivative intrinsic
bool IsDerivative(const std::string& name);
bool IsDerivative(ast::Intrinsic i);
/// Determines if the given |name| is a float classification intrinsic
/// @param name the name to check
/// @param i the intrinsic
/// @returns true if the given |name| is a float intrinsic
bool IsFloatClassificationIntrinsic(const std::string& name);
bool IsFloatClassificationIntrinsic(ast::Intrinsic i);
/// Determines if the given |name| is a texture operation intrinsic
/// @param name the name to check
/// @param i the intrinsic
/// @returns true if the given |name| is a texture operation intrinsic
bool IsTextureOperationIntrinsic(const std::string& name);
/// Determines if the given |name| is an intrinsic
/// @param name the name to check
/// @returns true if the given |name| is an intrinsic
bool IsIntrinsic(const std::string& name);
bool IsTextureIntrinsic(ast::Intrinsic i);
} // namespace intrinsic
} // namespace ast

17
src/ast/module.cc
View File

@ -27,15 +27,6 @@ Module::Module(Module&&) = default;
Module::~Module() = default;
Import* Module::FindImportByName(const std::string& name) const {
for (const auto& import : imports_) {
if (import->name() == name) {
return import.get();
}
}
return nullptr;
}
Function* Module::FindFunctionByName(const std::string& name) const {
for (const auto& func : functions_) {
if (func->name() == name) {
@ -56,11 +47,6 @@ Function* Module::FindFunctionByNameAndStage(const std::string& name,
}
bool Module::IsValid() const {
for (const auto& import : imports_) {
if (import == nullptr || !import->IsValid()) {
return false;
}
}
for (const auto& var : global_variables_) {
if (var == nullptr || !var->IsValid()) {
return false;
@ -84,9 +70,6 @@ std::string Module::to_str() const {
out << "Module{" << std::endl;
const auto indent = 2;
for (const auto& import : imports_) {
import->to_str(out, indent);
}
for (const auto& var : global_variables_) {
var->to_str(out, indent);
}

14
src/ast/module.h
View File

@ -21,7 +21,6 @@
#include <vector>
#include "src/ast/function.h"
#include "src/ast/import.h"
#include "src/ast/type/alias_type.h"
#include "src/ast/variable.h"
@ -36,18 +35,6 @@ class Module {
Module(Module&&);
~Module();
/// Add the given import to the module
/// @param import The import to add.
void AddImport(std::unique_ptr<Import> import) {
imports_.push_back(std::move(import));
}
/// @returns the imports for this module
const ImportList& imports() const { return imports_; }
/// Find the import of the given name
/// @param name The import name to search for
/// @returns the import with the given name if found, nullptr otherwise.
Import* FindImportByName(const std::string& name) const;
/// Add a global variable to the module
/// @param var the variable to add
void AddGlobalVariable(std::unique_ptr<Variable> var) {
@ -94,7 +81,6 @@ class Module {
private:
Module(const Module&) = delete;
ImportList imports_;
VariableList global_variables_;
// The alias types are owned by the type manager
std::vector<type::AliasType*> alias_types_;

59
src/ast/module_test.cc
View File

@ -19,7 +19,6 @@
#include "gmock/gmock.h"
#include "src/ast/function.h"
#include "src/ast/import.h"
#include "src/ast/type/f32_type.h"
#include "src/ast/variable.h"
@ -32,7 +31,7 @@ using ModuleTest = testing::Test;
TEST_F(ModuleTest, Creation) {
Module m;
EXPECT_EQ(m.imports().size(), 0u);
EXPECT_EQ(m.functions().size(), 0u);
}
TEST_F(ModuleTest, ToStrEmitsPreambleAndPostamble) {
@ -42,44 +41,6 @@ TEST_F(ModuleTest, ToStrEmitsPreambleAndPostamble) {
EXPECT_EQ(str, expected);
}
TEST_F(ModuleTest, Imports) {
Module m;
m.AddImport(std::make_unique<Import>("GLSL.std.430", "std::glsl"));
m.AddImport(std::make_unique<Import>("OpenCL.debug.100", "std::debug"));
EXPECT_EQ(2u, m.imports().size());
EXPECT_EQ("std::glsl", m.imports()[0]->name());
}
TEST_F(ModuleTest, ToStrWithImport) {
Module m;
m.AddImport(std::make_unique<Import>("GLSL.std.430", "std::glsl"));
const auto str = m.to_str();
EXPECT_EQ(str, R"(Module{
Import{"GLSL.std.430" as std::glsl}
}
)");
}
TEST_F(ModuleTest, LookupImport) {
Module m;
auto i = std::make_unique<Import>("GLSL.std.430", "std::glsl");
m.AddImport(std::move(i));
m.AddImport(std::make_unique<Import>("OpenCL.debug.100", "std::debug"));
auto* import = m.FindImportByName("std::glsl");
ASSERT_NE(nullptr, import);
EXPECT_EQ(import->path(), "GLSL.std.430");
EXPECT_EQ(import->name(), "std::glsl");
}
TEST_F(ModuleTest, LookupImportMissing) {
Module m;
EXPECT_EQ(nullptr, m.FindImportByName("Missing"));
}
TEST_F(ModuleTest, LookupFunction) {
type::F32Type f32;
Module m;
@ -100,24 +61,6 @@ TEST_F(ModuleTest, IsValid_Empty) {
EXPECT_TRUE(m.IsValid());
}
TEST_F(ModuleTest, IsValid_Import) {
Module m;
m.AddImport(std::make_unique<Import>("GLSL.std.430", "std::glsl"));
EXPECT_TRUE(m.IsValid());
}
TEST_F(ModuleTest, IsValid_Null_Import) {
Module m;
m.AddImport(nullptr);
EXPECT_FALSE(m.IsValid());
}
TEST_F(ModuleTest, IsValid_Invalid_Import) {
Module m;
m.AddImport(std::make_unique<Import>());
EXPECT_FALSE(m.IsValid());
}
TEST_F(ModuleTest, IsValid_GlobalVariable) {
type::F32Type f32;
auto var = std::make_unique<Variable>("var", StorageClass::kInput, &f32);

6
src/reader/spirv/function.cc
View File

@ -327,7 +327,7 @@ std::string GetGlslStd450FuncName(uint32_t ext_opcode) {
case GLSLstd450Normalize:
return "normalize";
case GLSLstd450FClamp:
return "fclamp";
return "clamp";
case GLSLstd450Length:
return "length";
default:
@ -2805,8 +2805,8 @@ TypedExpression FunctionEmitter::EmitGlslStd450ExtInst(
Fail() << "unhandled GLSL.std.450 instruction " << ext_opcode;
return {};
}
auto func = std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{parser_impl_.GlslStd450Prefix(), name});
auto func = std::make_unique<ast::IdentifierExpression>(name);
ast::ExpressionList operands;
// All parameters to GLSL.std.450 extended instructions are IDs.
for (uint32_t iarg = 2; iarg < inst.NumInOperands(); ++iarg) {

24
src/reader/spirv/function_glsl_std_450_test.cc
View File

@ -434,30 +434,28 @@ TEST_P(SpvParserTest_GlslStd450_Floating_FloatingFloatingFloating, Vector) {
INSTANTIATE_TEST_SUITE_P(Samples,
SpvParserTest_GlslStd450_Float_Floating,
::testing::Values(GlslStd450Case{
"Length", "std::glsl::length"}));
::testing::Values(GlslStd450Case{"Length", "length"}));
INSTANTIATE_TEST_SUITE_P(Samples,
SpvParserTest_GlslStd450_Float_FloatingFloating,
::testing::Values(GlslStd450Case{
"Distance", "std::glsl::distance"}));
::testing::Values(GlslStd450Case{"Distance",
"distance"}));
INSTANTIATE_TEST_SUITE_P(
Samples,
SpvParserTest_GlslStd450_Floating_Floating,
::testing::Values(GlslStd450Case{"Sin", "std::glsl::sin"},
GlslStd450Case{"Cos", "std::glsl::cos"},
GlslStd450Case{"Normalize", "std::glsl::normalize"}));
INSTANTIATE_TEST_SUITE_P(Samples,
SpvParserTest_GlslStd450_Floating_Floating,
::testing::Values(GlslStd450Case{"Sin", "sin"},
GlslStd450Case{"Cos", "cos"},
GlslStd450Case{"Normalize",
"normalize"}));
INSTANTIATE_TEST_SUITE_P(Samples,
SpvParserTest_GlslStd450_Floating_FloatingFloating,
::testing::Values(GlslStd450Case{"Atan2",
"std::glsl::atan2"}));
::testing::Values(GlslStd450Case{"Atan2", "atan2"}));
INSTANTIATE_TEST_SUITE_P(
Samples,
SpvParserTest_GlslStd450_Floating_FloatingFloatingFloating,
::testing::Values(GlslStd450Case{"FClamp", "std::glsl::fclamp"}));
::testing::Values(GlslStd450Case{"FClamp", "clamp"}));
} // namespace
} // namespace spirv

9
src/reader/spirv/parser_impl.cc
View File

@ -453,7 +453,6 @@ void ParserImpl::ResetInternalModule() {
type_mgr_ = nullptr;
deco_mgr_ = nullptr;
import_map_.clear();
glsl_std_450_imports_.clear();
}
@ -544,14 +543,6 @@ bool ParserImpl::RegisterExtendedInstructionImports() {
reinterpret_cast<const char*>(import.GetInOperand(0).words.data()));
// TODO(dneto): Handle other extended instruction sets when needed.
if (name == "GLSL.std.450") {
// Only create the AST import once, so we can use import name 'std::glsl'.
// This is a canonicalization.
if (glsl_std_450_imports_.empty()) {
auto ast_import =
std::make_unique<tint::ast::Import>(name, GlslStd450Prefix());
import_map_[import.result_id()] = ast_import.get();
ast_module_.AddImport(std::move(ast_import));
}
glsl_std_450_imports_.insert(import.result_id());
} else {
return Fail() << "Unrecognized extended instruction set: " << name;

6
src/reader/spirv/parser_impl.h
View File

@ -32,7 +32,6 @@
#include "source/opt/types.h"
#include "spirv-tools/libspirv.hpp"
#include "src/ast/expression.h"
#include "src/ast/import.h"
#include "src/ast/module.h"
#include "src/ast/struct_member_decoration.h"
#include "src/ast/type/alias_type.h"
@ -136,9 +135,6 @@ class ParserImpl : Reader {
return glsl_std_450_imports_;
}
/// @returns the import prefix to use for the GLSL.std.450 import.
std::string GlslStd450Prefix() const { return "std::glsl"; }
/// Converts a SPIR-V type to a Tint type, and saves it for fast lookup.
/// If the type is only used for builtins, then register that specially,
/// and return null.
@ -460,8 +456,6 @@ class ParserImpl : Reader {
// instruction is line 1.
std::unordered_map<const spvtools::opt::Instruction*, Source> inst_source_;
/// Maps a SPIR-V ID for an external instruction import to an AST import
std::unordered_map<uint32_t, ast::Import*> import_map_;
// The set of IDs that are imports of the GLSL.std.450 extended instruction
// sets.
std::unordered_set<uint32_t> glsl_std_450_imports_;

17
src/reader/spirv/parser_impl_import_test.cc
View File

@ -44,23 +44,6 @@ TEST_F(SpvParserTest, Import_ImportGlslStd450) {
EXPECT_TRUE(p->BuildAndParseInternalModule());
EXPECT_TRUE(p->error().empty());
EXPECT_THAT(p->glsl_std_450_imports(), ElementsAre(1));
const auto module_ast = p->module().to_str();
EXPECT_THAT(module_ast, HasSubstr(R"(Import{"GLSL.std.450" as std::glsl})"));
}
TEST_F(SpvParserTest, Import_ImportGlslStd450Twice) {
auto* p = parser(test::Assemble(R"(
%1 = OpExtInstImport "GLSL.std.450"
%42 = OpExtInstImport "GLSL.std.450"
)"));
EXPECT_TRUE(p->BuildAndParseInternalModule());
EXPECT_TRUE(p->error().empty());
EXPECT_THAT(p->glsl_std_450_imports(), UnorderedElementsAre(1, 42));
const auto module = p->module();
EXPECT_EQ(module.imports().size(), 1u);
const auto module_ast = module.to_str();
// TODO(dneto): Use a matcher to show there is only one import.
EXPECT_THAT(module_ast, HasSubstr(R"(Import{"GLSL.std.450" as std::glsl})"));
}
// TODO(dneto): We don't currently support other kinds of extended instruction

98
src/reader/wgsl/parser_impl.cc
View File

@ -203,7 +203,6 @@ void ParserImpl::translation_unit() {
// global_decl
// : SEMICOLON
// | import_decl SEMICOLON
// | global_variable_decl SEMICLON
// | global_constant_decl SEMICOLON
// | type_alias SEMICOLON
@ -218,19 +217,6 @@ void ParserImpl::global_decl() {
return;
}
auto import = import_decl();
if (has_error())
return;
if (import != nullptr) {
t = next();
if (!t.IsSemicolon()) {
set_error(t, "missing ';' for import");
return;
}
module_.AddImport(std::move(import));
return;
}
auto gv = global_variable_decl();
if (has_error())
return;
@ -281,69 +267,6 @@ void ParserImpl::global_decl() {
set_error(t);
}
// import_decl
// : IMPORT STRING_LITERAL AS (IDENT NAMESPACE)* IDENT
std::unique_ptr<ast::Import> ParserImpl::import_decl() {
auto t = peek();
if (!t.IsImport())
return {};
auto source = t.source();
next(); // consume the import token
t = next();
if (!t.IsStringLiteral()) {
set_error(t, "missing path for import");
return {};
}
auto path = t.to_str();
if (path.length() == 0) {
set_error(t, "import path must not be empty");
return {};
}
t = next();
if (!t.IsAs()) {
set_error(t, "missing 'as' for import");
return {};
}
std::string name = "";
for (;;) {
t = peek();
if (!t.IsIdentifier()) {
break;
}
next(); // consume the peek
name += t.to_str();
t = peek();
if (!t.IsNamespace()) {
break;
}
next(); // consume the peek
name += "::";
}
if (name.length() == 0) {
if (t.IsEof() || t.IsSemicolon()) {
set_error(t, "missing name for import");
} else {
set_error(t, "invalid name for import");
}
return {};
}
if (name.length() > 2) {
auto end = name.length() - 1;
if (name[end] == ':' && name[end - 1] == ':') {
set_error(t, "invalid name for import");
return {};
}
}
return std::make_unique<ast::Import>(source, path, name);
}
// global_variable_decl
// : variable_decoration_list variable_decl
// | variable_decoration_list variable_decl EQUAL const_expr
@ -2840,7 +2763,7 @@ std::unique_ptr<ast::BlockStatement> ParserImpl::continuing_stmt() {
}
// primary_expression
// : (IDENT NAMESPACE)* IDENT
// : IDENT
// | type_decl PAREN_LEFT argument_expression_list* PAREN_RIGHT
// | const_literal
// | paren_rhs_stmt
@ -2911,24 +2834,7 @@ std::unique_ptr<ast::Expression> ParserImpl::primary_expression() {
} else if (t.IsIdentifier()) {
next(); // Consume the peek
std::vector<std::string> ident;
ident.push_back(t.to_str());
for (;;) {
t = peek();
if (!t.IsNamespace())
break;
next(); // Consume the peek
t = next();
if (!t.IsIdentifier()) {
set_error(t, "identifier expected");
return nullptr;
}
ident.push_back(t.to_str());
}
return std::make_unique<ast::IdentifierExpression>(source,
std::move(ident));
return std::make_unique<ast::IdentifierExpression>(source, t.to_str());
}
auto* type = type_decl();

4
src/reader/wgsl/parser_impl.h
View File

@ -28,7 +28,6 @@
#include "src/ast/constructor_expression.h"
#include "src/ast/else_statement.h"
#include "src/ast/function.h"
#include "src/ast/import.h"
#include "src/ast/literal.h"
#include "src/ast/loop_statement.h"
#include "src/ast/module.h"
@ -123,9 +122,6 @@ class ParserImpl {
void translation_unit();
/// Parses the `global_decl` grammar element
void global_decl();
/// Parses the `import_decl grammar element
/// @returns the import object or nullptr if an error was encountered
std::unique_ptr<ast::Import> import_decl();
/// Parses a `global_variable_decl` grammar element
/// @returns the variable parsed or nullptr
std::unique_ptr<ast::Variable> global_variable_decl();

29
src/reader/wgsl/parser_impl_global_decl_test.cc
View File

@ -27,35 +27,6 @@ TEST_F(ParserImplTest, GlobalDecl_Semicolon) {
ASSERT_FALSE(p->has_error()) << p->error();
}
TEST_F(ParserImplTest, GlobalDecl_Import) {
auto* p = parser(R"(import "GLSL.std.430" as glsl;)");
p->global_decl();
ASSERT_FALSE(p->has_error()) << p->error();
auto m = p->module();
ASSERT_EQ(1u, m.imports().size());
const auto& import = m.imports()[0];
EXPECT_EQ("GLSL.std.430", import->path());
EXPECT_EQ("glsl", import->name());
}
TEST_F(ParserImplTest, GlobalDecl_Import_Invalid) {
auto* p = parser(R"(import as glsl;)");
p->global_decl();
ASSERT_TRUE(p->has_error());
EXPECT_EQ(p->error(), "1:8: missing path for import");
}
TEST_F(ParserImplTest, GlobalDecl_Import_Invalid_MissingSemicolon) {
auto* p = parser(R"(import "GLSL.std.430" as glsl)");
p->global_decl();
ASSERT_TRUE(p->has_error());
EXPECT_EQ(p->error(), "1:30: missing ';' for import");
}
TEST_F(ParserImplTest, GlobalDecl_GlobalVariable) {
auto* p = parser("var<out> a : vec2<i32> = vec2<i32>(1, 2);");
p->global_decl();

96
src/reader/wgsl/parser_impl_import_decl_test.cc
View File

@ -1,96 +0,0 @@
// Copyright 2020 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include "gtest/gtest.h"
#include "src/reader/wgsl/parser_impl.h"
#include "src/reader/wgsl/parser_impl_test_helper.h"
namespace tint {
namespace reader {
namespace wgsl {
namespace {
TEST_F(ParserImplTest, ImportDecl_Import) {
auto* p = parser(R"(import "GLSL.std.450" as glsl)");
auto import = p->import_decl();
ASSERT_NE(import, nullptr);
ASSERT_FALSE(p->has_error()) << p->error();
EXPECT_EQ("GLSL.std.450", import->path());
EXPECT_EQ("glsl", import->name());
EXPECT_EQ(1u, import->line());
EXPECT_EQ(1u, import->column());
}
TEST_F(ParserImplTest, ImportDecl_Import_WithNamespace) {
auto* p = parser(R"(import "GLSL.std.450" as std::glsl)");
auto import = p->import_decl();
ASSERT_NE(import, nullptr);
ASSERT_FALSE(p->has_error()) << p->error();
EXPECT_EQ("std::glsl", import->name());
}
TEST_F(ParserImplTest, ImportDecl_Invalid_MissingPath) {
auto* p = parser(R"(import as glsl)");
auto import = p->import_decl();
ASSERT_EQ(import, nullptr);
ASSERT_TRUE(p->has_error());
EXPECT_EQ(p->error(), "1:8: missing path for import");
}
TEST_F(ParserImplTest, ImportDecl_Invalid_EmptyPath) {
auto* p = parser(R"(import "" as glsl)");
auto import = p->import_decl();
ASSERT_EQ(import, nullptr);
ASSERT_TRUE(p->has_error());
EXPECT_EQ(p->error(), "1:8: import path must not be empty");
}
TEST_F(ParserImplTest, ImportDecl_Invalid_NameMissingTerminatingIdentifier) {
auto* p = parser(R"(import "GLSL.std.450" as glsl::)");
auto import = p->import_decl();
ASSERT_EQ(import, nullptr);
ASSERT_TRUE(p->has_error());
EXPECT_EQ(p->error(), "1:32: invalid name for import");
}
TEST_F(ParserImplTest, ImportDecl_Invalid_NameInvalid) {
auto* p = parser(R"(import "GLSL.std.450" as 12glsl)");
auto import = p->import_decl();
ASSERT_EQ(import, nullptr);
ASSERT_TRUE(p->has_error());
EXPECT_EQ(p->error(), "1:26: invalid name for import");
}
TEST_F(ParserImplTest, ImportDecl_Invalid_MissingName) {
auto* p = parser(R"(import "GLSL.std.450" as)");
auto import = p->import_decl();
ASSERT_EQ(import, nullptr);
ASSERT_TRUE(p->has_error());
EXPECT_EQ(p->error(), "1:25: missing name for import");
}
TEST_F(ParserImplTest, ImportDecl_Invalid_MissingAs) {
auto* p = parser(R"(import "GLSL.std.450" glsl)");
auto import = p->import_decl();
ASSERT_EQ(import, nullptr);
ASSERT_TRUE(p->has_error());
EXPECT_EQ(p->error(), "1:23: missing 'as' for import");
}
} // namespace
} // namespace wgsl
} // namespace reader
} // namespace tint

3
src/reader/wgsl/parser_impl_postfix_expression_test.cc
View File

@ -105,7 +105,7 @@ TEST_F(ParserImplTest, PostfixExpression_Call_Empty) {
}
TEST_F(ParserImplTest, PostfixExpression_Call_WithArgs) {
auto* p = parser("std::test(1, b, 2 + 3 / b)");
auto* p = parser("test(1, b, 2 + 3 / b)");
auto e = p->postfix_expression();
ASSERT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e, nullptr);
@ -115,7 +115,6 @@ TEST_F(ParserImplTest, PostfixExpression_Call_WithArgs) {
ASSERT_TRUE(c->func()->IsIdentifier());
auto* func = c->func()->AsIdentifier();
EXPECT_EQ(func->path(), "std");
EXPECT_EQ(func->name(), "test");
EXPECT_EQ(c->params().size(), 3u);

19
src/reader/wgsl/parser_impl_primary_expression_test.cc
View File

@ -43,25 +43,6 @@ TEST_F(ParserImplTest, PrimaryExpression_Ident) {
EXPECT_EQ(ident->name(), "a");
}
TEST_F(ParserImplTest, PrimaryExpression_Ident_WithNamespace) {
auto* p = parser("a::b::c::d");
auto e = p->primary_expression();
ASSERT_FALSE(p->has_error()) << p->error();
ASSERT_NE(e, nullptr);
ASSERT_TRUE(e->IsIdentifier());
auto* ident = e->AsIdentifier();
EXPECT_EQ(ident->path(), "a::b::c");
EXPECT_EQ(ident->name(), "d");
}
TEST_F(ParserImplTest, PrimaryExpression_Ident_MissingIdent) {
auto* p = parser("a::");
auto e = p->primary_expression();
ASSERT_TRUE(p->has_error());
ASSERT_EQ(e, nullptr);
EXPECT_EQ(p->error(), "1:4: identifier expected");
}
TEST_F(ParserImplTest, PrimaryExpression_TypeDecl) {
auto* p = parser("vec4<i32>(1, 2, 3, 4))");
auto e = p->primary_expression();

7
src/reader/wgsl/parser_impl_test.cc
View File

@ -30,8 +30,6 @@ TEST_F(ParserImplTest, Empty) {
TEST_F(ParserImplTest, Parses) {
auto* p = parser(R"(
import "GLSL.std.430" as glsl;
[[location(0)]] var<out> gl_FragColor : vec4<f32>;
[[stage(vertex)]]
@ -42,22 +40,19 @@ fn main() -> void {
ASSERT_TRUE(p->Parse()) << p->error();
auto m = p->module();
ASSERT_EQ(1u, m.imports().size());
ASSERT_EQ(1u, m.functions().size());
ASSERT_EQ(1u, m.global_variables().size());
}
TEST_F(ParserImplTest, HandlesError) {
auto* p = parser(R"(
import "GLSL.std.430" as glsl;
fn main() -> { # missing return type
return;
})");
ASSERT_FALSE(p->Parse());
ASSERT_TRUE(p->has_error());
EXPECT_EQ(p->error(), "4:15: unable to determine function return type");
EXPECT_EQ(p->error(), "2:15: unable to determine function return type");
}
TEST_F(ParserImplTest, GetRegisteredType) {

7
src/reader/wgsl/parser_test.cc
View File

@ -34,8 +34,6 @@ TEST_F(ParserTest, Parses) {
Context ctx;
Parser p(&ctx, R"(
import "GLSL.std.430" as glsl;
[[location(0)]] var<out> gl_FragColor : vec4<f32>;
[[stage(vertex)]]
@ -46,7 +44,6 @@ fn main() -> void {
ASSERT_TRUE(p.Parse()) << p.error();
auto m = p.module();
ASSERT_EQ(1u, m.imports().size());
ASSERT_EQ(1u, m.functions().size());
ASSERT_EQ(1u, m.global_variables().size());
}
@ -54,15 +51,13 @@ fn main() -> void {
TEST_F(ParserTest, HandlesError) {
Context ctx;
Parser p(&ctx, R"(
import "GLSL.std.430" as glsl;
fn main() -> { # missing return type
return;
})");
ASSERT_FALSE(p.Parse());
ASSERT_TRUE(p.has_error());
EXPECT_EQ(p.error(), "4:15: unable to determine function return type");
EXPECT_EQ(p.error(), "2:15: unable to determine function return type");
}
} // namespace

628
src/type_determiner.cc
View File

@ -17,7 +17,6 @@
#include <memory>
#include <vector>
#include "spirv/unified1/GLSL.std.450.h"
#include "src/ast/array_accessor_expression.h"
#include "src/ast/as_expression.h"
#include "src/ast/assignment_statement.h"
@ -56,104 +55,6 @@
#include "src/ast/variable_decl_statement.h"
namespace tint {
namespace {
// Most of these are floating-point general except the below which are only
// FP16 and FP32. We only have FP32 at this point so the below works, if we
// get FP64 support or otherwise we'll need to differentiate.
// * radians
// * degrees
// * sin, cos, tan
// * asin, acos, atan
// * sinh, cosh, tanh
// * asinh, acosh, atanh
// * exp, exp2
// * log, log2
enum class GlslDataType {
kFloatScalarOrVector,
kIntScalarOrVector,
kFloatVector,
kMatrix
};
struct GlslData {
const char* name;
uint8_t param_count;
uint32_t op_id;
GlslDataType type;
uint8_t vector_count;
};
constexpr const GlslData kGlslData[] = {
{"acos", 1, GLSLstd450Acos, GlslDataType::kFloatScalarOrVector, 0},
{"acosh", 1, GLSLstd450Acosh, GlslDataType::kFloatScalarOrVector, 0},
{"asin", 1, GLSLstd450Asin, GlslDataType::kFloatScalarOrVector, 0},
{"asinh", 1, GLSLstd450Asinh, GlslDataType::kFloatScalarOrVector, 0},
{"atan", 1, GLSLstd450Atan, GlslDataType::kFloatScalarOrVector, 0},
{"atan2", 2, GLSLstd450Atan2, GlslDataType::kFloatScalarOrVector, 0},
{"atanh", 1, GLSLstd450Atanh, GlslDataType::kFloatScalarOrVector, 0},
{"ceil", 1, GLSLstd450Ceil, GlslDataType::kFloatScalarOrVector, 0},
{"cos", 1, GLSLstd450Cos, GlslDataType::kFloatScalarOrVector, 0},
{"cosh", 1, GLSLstd450Cosh, GlslDataType::kFloatScalarOrVector, 0},
{"cross", 2, GLSLstd450Cross, GlslDataType::kFloatVector, 3},
{"degrees", 1, GLSLstd450Degrees, GlslDataType::kFloatScalarOrVector, 0},
{"determinant", 1, GLSLstd450Determinant, GlslDataType::kMatrix, 0},
{"distance", 2, GLSLstd450Distance, GlslDataType::kFloatScalarOrVector, 0},
{"exp", 1, GLSLstd450Exp, GlslDataType::kFloatScalarOrVector, 0},
{"exp2", 1, GLSLstd450Exp2, GlslDataType::kFloatScalarOrVector, 0},
{"fabs", 1, GLSLstd450FAbs, GlslDataType::kFloatScalarOrVector, 0},
{"faceforward", 3, GLSLstd450FaceForward,
GlslDataType::kFloatScalarOrVector, 0},
{"fclamp", 3, GLSLstd450FClamp, GlslDataType::kFloatScalarOrVector, 0},
{"findilsb", 1, GLSLstd450FindILsb, GlslDataType::kIntScalarOrVector, 0},
{"findumsb", 1, GLSLstd450FindUMsb, GlslDataType::kIntScalarOrVector, 0},
{"findsmsb", 1, GLSLstd450FindSMsb, GlslDataType::kIntScalarOrVector, 0},
{"floor", 1, GLSLstd450Floor, GlslDataType::kFloatScalarOrVector, 0},
{"fma", 3, GLSLstd450Fma, GlslDataType::kFloatScalarOrVector, 0},
{"fmax", 2, GLSLstd450FMax, GlslDataType::kFloatScalarOrVector, 0},
{"fmin", 2, GLSLstd450FMin, GlslDataType::kFloatScalarOrVector, 0},
{"fmix", 3, GLSLstd450FMix, GlslDataType::kFloatScalarOrVector, 0},
{"fract", 1, GLSLstd450Fract, GlslDataType::kFloatScalarOrVector, 0},
{"fsign", 1, GLSLstd450FSign, GlslDataType::kFloatScalarOrVector, 0},
{"interpolateatcentroid", 1, GLSLstd450InterpolateAtCentroid,
GlslDataType::kFloatScalarOrVector, 0},
{"inversesqrt", 1, GLSLstd450InverseSqrt,
GlslDataType::kFloatScalarOrVector, 0},
{"length", 1, GLSLstd450Length, GlslDataType::kFloatScalarOrVector, 0},
{"log", 1, GLSLstd450Log, GlslDataType::kFloatScalarOrVector, 0},
{"log2", 1, GLSLstd450Log2, GlslDataType::kFloatScalarOrVector, 0},
{"matrixinverse", 1, GLSLstd450MatrixInverse, GlslDataType::kMatrix, 0},
{"nclamp", 3, GLSLstd450NClamp, GlslDataType::kFloatScalarOrVector, 0},
{"nmax", 2, GLSLstd450NMax, GlslDataType::kFloatScalarOrVector, 0},
{"nmin", 2, GLSLstd450NMin, GlslDataType::kFloatScalarOrVector, 0},
{"normalize", 1, GLSLstd450Normalize, GlslDataType::kFloatScalarOrVector,
0},
{"pow", 2, GLSLstd450Pow, GlslDataType::kFloatScalarOrVector, 0},
{"radians", 1, GLSLstd450Radians, GlslDataType::kFloatScalarOrVector, 0},
{"reflect", 2, GLSLstd450Reflect, GlslDataType::kFloatScalarOrVector, 0},
{"round", 1, GLSLstd450Round, GlslDataType::kFloatScalarOrVector, 0},
{"roundeven", 1, GLSLstd450RoundEven, GlslDataType::kFloatScalarOrVector,
0},
{"sabs", 1, GLSLstd450SAbs, GlslDataType::kIntScalarOrVector, 0},
{"sclamp", 3, GLSLstd450SClamp, GlslDataType::kIntScalarOrVector, 0},
{"sin", 1, GLSLstd450Sin, GlslDataType::kFloatScalarOrVector, 0},
{"sinh", 1, GLSLstd450Sinh, GlslDataType::kFloatScalarOrVector, 0},
{"smax", 2, GLSLstd450SMax, GlslDataType::kIntScalarOrVector, 0},
{"smin", 2, GLSLstd450SMin, GlslDataType::kIntScalarOrVector, 0},
{"smoothstep", 3, GLSLstd450SmoothStep, GlslDataType::kFloatScalarOrVector,
0},
{"sqrt", 1, GLSLstd450Sqrt, GlslDataType::kFloatScalarOrVector, 0},
{"ssign", 1, GLSLstd450SSign, GlslDataType::kIntScalarOrVector, 0},
{"step", 2, GLSLstd450Step, GlslDataType::kFloatScalarOrVector, 0},
{"tan", 1, GLSLstd450Tan, GlslDataType::kFloatScalarOrVector, 0},
{"tanh", 1, GLSLstd450Tanh, GlslDataType::kFloatScalarOrVector, 0},
{"trunc", 1, GLSLstd450Trunc, GlslDataType::kFloatScalarOrVector, 0},
{"uclamp", 3, GLSLstd450UClamp, GlslDataType::kIntScalarOrVector, 0},
{"umax", 2, GLSLstd450UMax, GlslDataType::kIntScalarOrVector, 0},
{"umin", 2, GLSLstd450UMin, GlslDataType::kIntScalarOrVector, 0},
};
constexpr const uint32_t kGlslDataCount = sizeof(kGlslData) / sizeof(GlslData);
} // namespace
TypeDeterminer::TypeDeterminer(Context* ctx, ast::Module* mod)
: ctx_(*ctx), mod_(mod) {}
@ -471,6 +372,9 @@ bool TypeDeterminer::DetermineAs(ast::AsExpression* expr) {
}
bool TypeDeterminer::DetermineCall(ast::CallExpression* expr) {
if (!DetermineResultType(expr->func())) {
return false;
}
if (!DetermineResultType(expr->params())) {
return false;
}
@ -481,31 +385,10 @@ bool TypeDeterminer::DetermineCall(ast::CallExpression* expr) {
if (expr->func()->IsIdentifier()) {
auto* ident = expr->func()->AsIdentifier();
if (ast::intrinsic::IsIntrinsic(ident->name())) {
if (!DetermineIntrinsic(ident->name(), expr))
return false;
} else if (ident->has_path()) {
auto* imp = mod_->FindImportByName(ident->path());
if (imp == nullptr) {
set_error(expr->source(), "Unable to find import for " + ident->name());
if (ident->IsIntrinsic()) {
if (!DetermineIntrinsic(ident, expr)) {
return false;
}
uint32_t ext_id = 0;
auto* result_type = GetImportData(expr->source(), imp->path(),
ident->name(), expr->params(), &ext_id);
if (result_type == nullptr) {
if (error_.empty()) {
set_error(expr->source(),
"Unable to determine result type for GLSL expression " +
ident->name());
}
return false;
}
imp->AddMethodId(ident->name(), ext_id);
expr->func()->set_result_type(result_type);
} else {
if (current_function_) {
caller_to_callee_[current_function_->name()].push_back(ident->name());
@ -547,41 +430,111 @@ bool TypeDeterminer::DetermineCall(ast::CallExpression* expr) {
return true;
}
bool TypeDeterminer::DetermineIntrinsic(const std::string& name,
namespace {
enum class IntrinsicDataType {
kFloatOrIntScalarOrVector,
kFloatScalarOrVector,
kIntScalarOrVector,
kFloatVector,
kMatrix,
};
struct IntrinsicData {
ast::Intrinsic intrinsic;
uint8_t param_count;
IntrinsicDataType data_type;
uint8_t vector_size;
};
// Note, this isn't all the intrinsics. Some are handled specially before
// we get to the generic code. See the DetermineIntrinsic code below.
constexpr const IntrinsicData kIntrinsicData[] = {
{ast::Intrinsic::kAbs, 1, IntrinsicDataType::kFloatOrIntScalarOrVector, 0},
{ast::Intrinsic::kAcos, 1, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kAsin, 1, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kAtan, 1, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kAtan2, 2, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kCeil, 1, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kClamp, 3, IntrinsicDataType::kFloatOrIntScalarOrVector,
0},
{ast::Intrinsic::kCos, 1, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kCosh, 1, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kCountOneBits, 1, IntrinsicDataType::kIntScalarOrVector,
0},
{ast::Intrinsic::kCross, 2, IntrinsicDataType::kFloatVector, 3},
{ast::Intrinsic::kDeterminant, 1, IntrinsicDataType::kMatrix, 0},
{ast::Intrinsic::kDistance, 2, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kExp, 1, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kExp2, 1, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kFaceForward, 3, IntrinsicDataType::kFloatScalarOrVector,
0},
{ast::Intrinsic::kFloor, 1, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kFma, 3, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kFract, 1, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kFrexp, 2, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kInverseSqrt, 1, IntrinsicDataType::kFloatScalarOrVector,
0},
{ast::Intrinsic::kLdexp, 2, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kLength, 1, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kLog, 1, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kLog2, 1, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kMax, 2, IntrinsicDataType::kFloatOrIntScalarOrVector, 0},
{ast::Intrinsic::kMin, 2, IntrinsicDataType::kFloatOrIntScalarOrVector, 0},
{ast::Intrinsic::kMix, 3, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kModf, 2, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kNormalize, 1, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kPow, 2, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kReflect, 2, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kReverseBits, 1, IntrinsicDataType::kIntScalarOrVector, 0},
{ast::Intrinsic::kRound, 1, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kSign, 1, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kSin, 1, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kSinh, 1, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kSmoothStep, 3, IntrinsicDataType::kFloatScalarOrVector,
0},
{ast::Intrinsic::kSqrt, 1, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kStep, 2, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kTan, 1, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kTanh, 1, IntrinsicDataType::kFloatScalarOrVector, 0},
{ast::Intrinsic::kTrunc, 1, IntrinsicDataType::kFloatScalarOrVector, 0},
};
constexpr const uint32_t kIntrinsicDataCount =
sizeof(kIntrinsicData) / sizeof(IntrinsicData);
} // namespace
bool TypeDeterminer::DetermineIntrinsic(ast::IdentifierExpression* ident,
ast::CallExpression* expr) {
if (ast::intrinsic::IsDerivative(name)) {
if (ast::intrinsic::IsDerivative(ident->intrinsic())) {
if (expr->params().size() != 1) {
set_error(expr->source(), "incorrect number of parameters for " + name);
set_error(expr->source(),
"incorrect number of parameters for " + ident->name());
return false;
}
// The result type must be the same as the type of the parameter.
auto& param = expr->params()[0];
if (!DetermineResultType(param.get())) {
return false;
}
expr->func()->set_result_type(param->result_type()->UnwrapPtrIfNeeded());
auto* param_type = expr->params()[0]->result_type()->UnwrapPtrIfNeeded();
expr->func()->set_result_type(param_type);
return true;
}
if (name == "any" || name == "all") {
if (ident->intrinsic() == ast::Intrinsic::kAny ||
ident->intrinsic() == ast::Intrinsic::kAll) {
expr->func()->set_result_type(
ctx_.type_mgr().Get(std::make_unique<ast::type::BoolType>()));
return true;
}
if (ast::intrinsic::IsFloatClassificationIntrinsic(name)) {
if (ast::intrinsic::IsFloatClassificationIntrinsic(ident->intrinsic())) {
if (expr->params().size() != 1) {
set_error(expr->source(), "incorrect number of parameters for " + name);
set_error(expr->source(),
"incorrect number of parameters for " + ident->name());
return false;
}
auto* bool_type =
ctx_.type_mgr().Get(std::make_unique<ast::type::BoolType>());
auto& param = expr->params()[0];
if (!DetermineResultType(param.get())) {
return false;
}
auto* param_type = param->result_type()->UnwrapPtrIfNeeded();
auto* param_type = expr->params()[0]->result_type()->UnwrapPtrIfNeeded();
if (param_type->IsVector()) {
expr->func()->set_result_type(
ctx_.type_mgr().Get(std::make_unique<ast::type::VectorType>(
@ -591,31 +544,31 @@ bool TypeDeterminer::DetermineIntrinsic(const std::string& name,
}
return true;
}
if (ast::intrinsic::IsTextureOperationIntrinsic(name)) {
if (ast::intrinsic::IsTextureIntrinsic(ident->intrinsic())) {
// TODO: Remove the LOD param from textureLoad on storage textures when
// https://github.com/gpuweb/gpuweb/pull/1032 gets merged.
uint32_t num_of_params =
(name == "textureLoad" || name == "textureSample") ? 3 : 4;
(ident->intrinsic() == ast::Intrinsic::kTextureLoad ||
ident->intrinsic() == ast::Intrinsic::kTextureSample)
? 3
: 4;
if (expr->params().size() != num_of_params) {
set_error(expr->source(),
"incorrect number of parameters for " + name + ", got " +
std::to_string(expr->params().size()) + " and expected " +
std::to_string(num_of_params));
"incorrect number of parameters for " + ident->name() +
", got " + std::to_string(expr->params().size()) +
" and expected " + std::to_string(num_of_params));
return false;
}
if (name == "textureSampleCompare") {
if (ident->intrinsic() == ast::Intrinsic::kTextureSampleCompare) {
expr->func()->set_result_type(
ctx_.type_mgr().Get(std::make_unique<ast::type::F32Type>()));
return true;
}
auto& texture_param = expr->params()[0];
if (!DetermineResultType(texture_param.get())) {
return false;
}
if (!texture_param->result_type()->UnwrapPtrIfNeeded()->IsTexture()) {
set_error(expr->source(), "invalid first argument for " + name);
set_error(expr->source(), "invalid first argument for " + ident->name());
return false;
}
ast::type::TextureType* texture =
@ -623,7 +576,7 @@ bool TypeDeterminer::DetermineIntrinsic(const std::string& name,
if (!texture->IsStorage() &&
!(texture->IsSampled() || texture->IsMultisampled())) {
set_error(expr->source(), "invalid texture for " + name);
set_error(expr->source(), "invalid texture for " + ident->name());
return false;
}
@ -642,29 +595,22 @@ bool TypeDeterminer::DetermineIntrinsic(const std::string& name,
ctx_.type_mgr().Get(std::make_unique<ast::type::VectorType>(type, 4)));
return true;
}
if (name == "dot") {
if (ident->intrinsic() == ast::Intrinsic::kDot) {
expr->func()->set_result_type(
ctx_.type_mgr().Get(std::make_unique<ast::type::F32Type>()));
return true;
}
if (name == "outerProduct") {
if (ident->intrinsic() == ast::Intrinsic::kOuterProduct) {
if (expr->params().size() != 2) {
set_error(expr->source(),
"incorrect number of parameters for outer_product");
"incorrect number of parameters for " + ident->name());
return false;
}
auto& param0 = expr->params()[0];
auto& param1 = expr->params()[1];
if (!DetermineResultType(param0.get()) ||
!DetermineResultType(param1.get())) {
return false;
}
auto* param0_type = param0->result_type()->UnwrapPtrIfNeeded();
auto* param1_type = param1->result_type()->UnwrapPtrIfNeeded();
auto* param0_type = expr->params()[0]->result_type()->UnwrapPtrIfNeeded();
auto* param1_type = expr->params()[1]->result_type()->UnwrapPtrIfNeeded();
if (!param0_type->IsVector() || !param1_type->IsVector()) {
set_error(expr->source(), "invalid parameter type for outer_product");
set_error(expr->source(), "invalid parameter type for " + ident->name());
return false;
}
@ -674,24 +620,121 @@ bool TypeDeterminer::DetermineIntrinsic(const std::string& name,
param0_type->AsVector()->size(), param1_type->AsVector()->size())));
return true;
}
if (name == "select") {
if (ident->intrinsic() == ast::Intrinsic::kSelect) {
if (expr->params().size() != 3) {
set_error(expr->source(),
"incorrect number of parameters for select expected 3 got " +
std::to_string(expr->params().size()));
set_error(expr->source(), "incorrect number of parameters for " +
ident->name() + " expected 3 got " +
std::to_string(expr->params().size()));
return false;
}
// The result type must be the same as the type of the parameter.
auto& param = expr->params()[0];
if (!DetermineResultType(param.get())) {
return false;
}
expr->func()->set_result_type(param->result_type()->UnwrapPtrIfNeeded());
auto* param_type = expr->params()[0]->result_type()->UnwrapPtrIfNeeded();
expr->func()->set_result_type(param_type);
return true;
}
return false;
const IntrinsicData* data = nullptr;
for (uint32_t i = 0; i < kIntrinsicDataCount; ++i) {
if (ident->intrinsic() == kIntrinsicData[i].intrinsic) {
data = &kIntrinsicData[i];
break;
}
}
if (data == nullptr) {
return false;
}
if (expr->params().size() != data->param_count) {
set_error(expr->source(), "incorrect number of parameters for " +
ident->name() + ". Expected " +
std::to_string(data->param_count) + " got " +
std::to_string(expr->params().size()));
return false;
}
std::vector<ast::type::Type*> result_types;
for (uint32_t i = 0; i < data->param_count; ++i) {
result_types.push_back(
expr->params()[i]->result_type()->UnwrapPtrIfNeeded());
switch (data->data_type) {
case IntrinsicDataType::kFloatOrIntScalarOrVector:
if (!result_types.back()->is_float_scalar_or_vector() &&
!result_types.back()->is_integer_scalar_or_vector()) {
set_error(expr->source(),
"incorrect type for " + ident->name() + ". " +
"Requires float or int, scalar or vector values");
return false;
}
break;
case IntrinsicDataType::kFloatScalarOrVector:
if (!result_types.back()->is_float_scalar_or_vector()) {
set_error(expr->source(),
"incorrect type for " + ident->name() + ". " +
"Requires float scalar or float vector values");
return false;
}
break;
case IntrinsicDataType::kIntScalarOrVector:
if (!result_types.back()->is_integer_scalar_or_vector()) {
set_error(expr->source(),
"incorrect type for " + ident->name() + ". " +
"Requires integer scalar or integer vector values");
return false;
}
break;
case IntrinsicDataType::kFloatVector:
if (!result_types.back()->is_float_vector()) {
set_error(expr->source(), "incorrect type for " + ident->name() +
". " + "Requires float vector values");
return false;
}
if (data->vector_size > 0 &&
result_types.back()->AsVector()->size() != data->vector_size) {
set_error(expr->source(), "incorrect vector size for " +
ident->name() + ". " + "Requires " +
std::to_string(data->vector_size) +
" elements");
return false;
}
break;
case IntrinsicDataType::kMatrix:
if (!result_types.back()->IsMatrix()) {
set_error(expr->source(), "incorrect type for " + ident->name() +
". Requires matrix value");
return false;
}
break;
}
}
// Verify all the parameter types match
for (size_t i = 1; i < data->param_count; ++i) {
if (result_types[0] != result_types[i]) {
set_error(expr->source(),
"mismatched parameter types for " + ident->name());
return false;
}
}
// Handle functions which aways return the type, even if a vector is
// provided.
if (ident->intrinsic() == ast::Intrinsic::kLength ||
ident->intrinsic() == ast::Intrinsic::kDistance) {
expr->func()->set_result_type(result_types[0]->is_float_scalar()
? result_types[0]
: result_types[0]->AsVector()->type());
return true;
}
// The determinant returns the component type of the columns
if (ident->intrinsic() == ast::Intrinsic::kDeterminant) {
expr->func()->set_result_type(result_types[0]->AsMatrix()->type());
return true;
}
expr->func()->set_result_type(result_types[0]);
return true;
}
bool TypeDeterminer::DetermineCast(ast::CastExpression* expr) {
@ -719,12 +762,6 @@ bool TypeDeterminer::DetermineConstructor(ast::ConstructorExpression* expr) {
}
bool TypeDeterminer::DetermineIdentifier(ast::IdentifierExpression* expr) {
if (expr->has_path()) {
set_error(expr->source(),
"determine identifier should not be called with imports");
return false;
}
auto name = expr->name();
ast::Variable* var;
if (variable_stack_.get(name, &var)) {
@ -750,9 +787,147 @@ bool TypeDeterminer::DetermineIdentifier(ast::IdentifierExpression* expr) {
return true;
}
SetIntrinsicIfNeeded(expr);
return true;
}
void TypeDeterminer::SetIntrinsicIfNeeded(ast::IdentifierExpression* ident) {
if (ident->name() == "abs") {
ident->set_intrinsic(ast::Intrinsic::kAbs);
} else if (ident->name() == "acos") {
ident->set_intrinsic(ast::Intrinsic::kAcos);
} else if (ident->name() == "all") {
ident->set_intrinsic(ast::Intrinsic::kAll);
} else if (ident->name() == "any") {
ident->set_intrinsic(ast::Intrinsic::kAny);
} else if (ident->name() == "asin") {
ident->set_intrinsic(ast::Intrinsic::kAsin);
} else if (ident->name() == "atan") {
ident->set_intrinsic(ast::Intrinsic::kAtan);
} else if (ident->name() == "atan2") {
ident->set_intrinsic(ast::Intrinsic::kAtan2);
} else if (ident->name() == "ceil") {
ident->set_intrinsic(ast::Intrinsic::kCeil);
} else if (ident->name() == "clamp") {
ident->set_intrinsic(ast::Intrinsic::kClamp);
} else if (ident->name() == "cos") {
ident->set_intrinsic(ast::Intrinsic::kCos);
} else if (ident->name() == "cosh") {
ident->set_intrinsic(ast::Intrinsic::kCosh);
} else if (ident->name() == "countOneBits") {
ident->set_intrinsic(ast::Intrinsic::kCountOneBits);
} else if (ident->name() == "cross") {
ident->set_intrinsic(ast::Intrinsic::kCross);
} else if (ident->name() == "determinant") {
ident->set_intrinsic(ast::Intrinsic::kDeterminant);
} else if (ident->name() == "distance") {
ident->set_intrinsic(ast::Intrinsic::kDistance);
} else if (ident->name() == "dot") {
ident->set_intrinsic(ast::Intrinsic::kDot);
} else if (ident->name() == "dpdx") {
ident->set_intrinsic(ast::Intrinsic::kDpdx);
} else if (ident->name() == "dpdxCoarse") {
ident->set_intrinsic(ast::Intrinsic::kDpdxCoarse);
} else if (ident->name() == "dpdxFine") {
ident->set_intrinsic(ast::Intrinsic::kDpdxFine);
} else if (ident->name() == "dpdy") {
ident->set_intrinsic(ast::Intrinsic::kDpdy);
} else if (ident->name() == "dpdyCoarse") {
ident->set_intrinsic(ast::Intrinsic::kDpdyCoarse);
} else if (ident->name() == "dpdyFine") {
ident->set_intrinsic(ast::Intrinsic::kDpdyFine);
} else if (ident->name() == "exp") {
ident->set_intrinsic(ast::Intrinsic::kExp);
} else if (ident->name() == "exp2") {
ident->set_intrinsic(ast::Intrinsic::kExp2);
} else if (ident->name() == "faceForward") {
ident->set_intrinsic(ast::Intrinsic::kFaceForward);
} else if (ident->name() == "floor") {
ident->set_intrinsic(ast::Intrinsic::kFloor);
} else if (ident->name() == "fma") {
ident->set_intrinsic(ast::Intrinsic::kFma);
} else if (ident->name() == "fract") {
ident->set_intrinsic(ast::Intrinsic::kFract);
} else if (ident->name() == "frexp") {
ident->set_intrinsic(ast::Intrinsic::kFrexp);
} else if (ident->name() == "fwidth") {
ident->set_intrinsic(ast::Intrinsic::kFwidth);
} else if (ident->name() == "fwidthCoarse") {
ident->set_intrinsic(ast::Intrinsic::kFwidthCoarse);
} else if (ident->name() == "fwidthFine") {
ident->set_intrinsic(ast::Intrinsic::kFwidthFine);
} else if (ident->name() == "inverseSqrt") {
ident->set_intrinsic(ast::Intrinsic::kInverseSqrt);
} else if (ident->name() == "isFinite") {
ident->set_intrinsic(ast::Intrinsic::kIsFinite);
} else if (ident->name() == "isInf") {
ident->set_intrinsic(ast::Intrinsic::kIsInf);
} else if (ident->name() == "isNan") {
ident->set_intrinsic(ast::Intrinsic::kIsNan);
} else if (ident->name() == "isNormal") {
ident->set_intrinsic(ast::Intrinsic::kIsNormal);
} else if (ident->name() == "ldexp") {
ident->set_intrinsic(ast::Intrinsic::kLdexp);
} else if (ident->name() == "length") {
ident->set_intrinsic(ast::Intrinsic::kLength);
} else if (ident->name() == "log") {
ident->set_intrinsic(ast::Intrinsic::kLog);
} else if (ident->name() == "log2") {
ident->set_intrinsic(ast::Intrinsic::kLog2);
} else if (ident->name() == "max") {
ident->set_intrinsic(ast::Intrinsic::kMax);
} else if (ident->name() == "min") {
ident->set_intrinsic(ast::Intrinsic::kMin);
} else if (ident->name() == "mix") {
ident->set_intrinsic(ast::Intrinsic::kMix);
} else if (ident->name() == "modf") {
ident->set_intrinsic(ast::Intrinsic::kModf);
} else if (ident->name() == "normalize") {
ident->set_intrinsic(ast::Intrinsic::kNormalize);
} else if (ident->name() == "outerProduct") {
ident->set_intrinsic(ast::Intrinsic::kOuterProduct);
} else if (ident->name() == "pow") {
ident->set_intrinsic(ast::Intrinsic::kPow);
} else if (ident->name() == "reflect") {
ident->set_intrinsic(ast::Intrinsic::kReflect);
} else if (ident->name() == "reverseBits") {
ident->set_intrinsic(ast::Intrinsic::kReverseBits);
} else if (ident->name() == "round") {
ident->set_intrinsic(ast::Intrinsic::kRound);
} else if (ident->name() == "select") {
ident->set_intrinsic(ast::Intrinsic::kSelect);
} else if (ident->name() == "sign") {
ident->set_intrinsic(ast::Intrinsic::kSign);
} else if (ident->name() == "sin") {
ident->set_intrinsic(ast::Intrinsic::kSin);
} else if (ident->name() == "sinh") {
ident->set_intrinsic(ast::Intrinsic::kSinh);
} else if (ident->name() == "smoothStep") {
ident->set_intrinsic(ast::Intrinsic::kSmoothStep);
} else if (ident->name() == "sqrt") {
ident->set_intrinsic(ast::Intrinsic::kSqrt);
} else if (ident->name() == "step") {
ident->set_intrinsic(ast::Intrinsic::kStep);
} else if (ident->name() == "tan") {
ident->set_intrinsic(ast::Intrinsic::kTan);
} else if (ident->name() == "tanh") {
ident->set_intrinsic(ast::Intrinsic::kTanh);
} else if (ident->name() == "textureLoad") {
ident->set_intrinsic(ast::Intrinsic::kTextureLoad);
} else if (ident->name() == "textureSample") {
ident->set_intrinsic(ast::Intrinsic::kTextureSample);
} else if (ident->name() == "textureSampleBias") {
ident->set_intrinsic(ast::Intrinsic::kTextureSampleBias);
} else if (ident->name() == "textureSampleCompare") {
ident->set_intrinsic(ast::Intrinsic::kTextureSampleCompare);
} else if (ident->name() == "textureSampleLevel") {
ident->set_intrinsic(ast::Intrinsic::kTextureSampleLevel);
} else if (ident->name() == "trunc") {
ident->set_intrinsic(ast::Intrinsic::kTrunc);
}
}
bool TypeDeterminer::DetermineMemberAccessor(
ast::MemberAccessorExpression* expr) {
if (!DetermineResultType(expr->structure())) {
@ -962,101 +1137,4 @@ bool TypeDeterminer::DetermineStorageTextureSubtype(
return false;
}
ast::type::Type* TypeDeterminer::GetImportData(
const Source& source,
const std::string& path,
const std::string& name,
const ast::ExpressionList& params,
uint32_t* id) {
if (path != "GLSL.std.450") {
set_error(source, "unknown import path " + path);
return nullptr;
}
const GlslData* data = nullptr;
for (uint32_t i = 0; i < kGlslDataCount; ++i) {
if (name == kGlslData[i].name) {
data = &kGlslData[i];
break;
}
}
if (data == nullptr) {
return nullptr;
}
if (params.size() != data->param_count) {
set_error(source, "incorrect number of parameters for " + name +
". Expected " + std::to_string(data->param_count) +
" got " + std::to_string(params.size()));
return nullptr;
}
std::vector<ast::type::Type*> result_types;
for (uint32_t i = 0; i < data->param_count; ++i) {
result_types.push_back(params[i]->result_type()->UnwrapPtrIfNeeded());
switch (data->type) {
case GlslDataType::kFloatScalarOrVector:
if (!result_types.back()->is_float_scalar_or_vector()) {
set_error(source, "incorrect type for " + name + ". " +
"Requires float scalar or float vector values");
return nullptr;
}
break;
case GlslDataType::kIntScalarOrVector:
if (!result_types.back()->is_integer_scalar_or_vector()) {
set_error(source,
"incorrect type for " + name + ". " +
"Requires integer scalar or integer vector values");
return nullptr;
}
break;
case GlslDataType::kFloatVector:
if (!result_types.back()->is_float_vector()) {
set_error(source, "incorrect type for " + name + ". " +
"Requires float vector values");
return nullptr;
}
if (data->vector_count > 0 &&
result_types.back()->AsVector()->size() != data->vector_count) {
set_error(source,
"incorrect vector size for " + name + ". " + "Requires " +
std::to_string(data->vector_count) + " elements");
return nullptr;
}
break;
case GlslDataType::kMatrix:
if (!result_types.back()->IsMatrix()) {
set_error(source,
"incorrect type for " + name + ". Requires matrix value");
return nullptr;
}
break;
}
}
// Verify all the parameter types match
for (size_t i = 1; i < data->param_count; ++i) {
if (result_types[0] != result_types[i]) {
error_ = "mismatched parameter types for " + name;
return nullptr;
}
}
*id = data->op_id;
// Handle functions which aways return the type, even if a vector is provided.
if (name == "length" || name == "distance") {
return result_types[0]->is_float_scalar()
? result_types[0]
: result_types[0]->AsVector()->type();
}
// The determinant returns the component type of the columns
if (name == "determinant") {
return result_types[0]->AsMatrix()->type();
}
return result_types[0];
}
} // namespace tint

7
src/type_determiner.h
View File

@ -107,6 +107,10 @@ class TypeDeterminer {
const ast::ExpressionList& params,
uint32_t* id);
/// Sets the intrinsic data information for the identifier if needed
/// @param ident the identifier expression
void SetIntrinsicIfNeeded(ast::IdentifierExpression* ident);
private:
void set_error(const Source& src, const std::string& msg);
void set_referenced_from_function_if_needed(ast::Variable* var);
@ -119,7 +123,8 @@ class TypeDeterminer {
bool DetermineCast(ast::CastExpression* expr);
bool DetermineConstructor(ast::ConstructorExpression* expr);
bool DetermineIdentifier(ast::IdentifierExpression* expr);
bool DetermineIntrinsic(const std::string& name, ast::CallExpression* expr);
bool DetermineIntrinsic(ast::IdentifierExpression* name,
ast::CallExpression* expr);
bool DetermineMemberAccessor(ast::MemberAccessorExpression* expr);
bool DetermineUnaryOp(ast::UnaryOpExpression* expr);

1841
src/type_determiner_test.cc
View File

File diff suppressed because it is too large Load Diff

17
src/validator_impl.cc
View File

@ -46,9 +46,6 @@ bool ValidatorImpl::Validate(const ast::Module* module) {
if (!ValidateGlobalVariables(module->global_variables())) {
return false;
}
if (!CheckImports(module)) {
return false;
}
if (!ValidateFunctions(module->functions())) {
return false;
}
@ -330,9 +327,7 @@ bool ValidatorImpl::ValidateCallExpr(const ast::CallExpression* expr) {
if (expr->func()->IsIdentifier()) {
auto* ident = expr->func()->AsIdentifier();
auto func_name = ident->name();
if (ident->has_path()) {
// TODO(sarahM0): validate import statements
} else if (ast::intrinsic::IsIntrinsic(ident->name())) {
if (ident->IsIntrinsic()) {
// TODO(sarahM0): validate intrinsics - tied with type-determiner
} else {
if (!function_stack_.has(func_name)) {
@ -433,14 +428,4 @@ bool ValidatorImpl::ValidateIdentifier(const ast::IdentifierExpression* ident) {
return true;
}
bool ValidatorImpl::CheckImports(const ast::Module* module) {
for (const auto& import : module->imports()) {
if (import->path() != "GLSL.std.450") {
set_error(import->source(), "v-0001: unknown import: " + import->path());
return false;
}
}
return true;
}
} // namespace tint

4
src/validator_impl.h
View File

@ -76,10 +76,6 @@ class ValidatorImpl {
/// @param assign the assignment to check
/// @returns true if the validation was successful
bool ValidateAssign(const ast::AssignmentStatement* assign);
/// Validates v-0001: Only allowed import is "GLSL.std.450"
/// @param module the modele to check imports
/// @returns ture if input complies with v-0001 rule
bool CheckImports(const ast::Module* module);
/// Validates an expression
/// @param expr the expression to check
/// @return true if the expression is valid

27
src/validator_test.cc
View File

@ -63,33 +63,6 @@ namespace {
class ValidatorTest : public ValidatorTestHelper, public testing::Test {};
TEST_F(ValidatorTest, Import) {
ast::Module m;
m.AddImport(std::make_unique<ast::Import>("GLSL.std.450", "glsl"));
EXPECT_TRUE(v()->CheckImports(&m));
}
TEST_F(ValidatorTest, Import_Fail_NotGLSL) {
ast::Module m;
m.AddImport(
std::make_unique<ast::Import>(Source{12, 34}, "not.GLSL", "glsl"));
EXPECT_FALSE(v()->CheckImports(&m));
ASSERT_TRUE(v()->has_error());
EXPECT_EQ(v()->error(), "12:34: v-0001: unknown import: not.GLSL");
}
TEST_F(ValidatorTest, Import_Fail_Typo) {
ast::Module m;
m.AddImport(
std::make_unique<ast::Import>(Source{12, 34}, "GLSL.std.4501", "glsl"));
EXPECT_FALSE(v()->CheckImports(&m));
ASSERT_TRUE(v()->has_error());
EXPECT_EQ(v()->error(), "12:34: v-0001: unknown import: GLSL.std.4501");
}
TEST_F(ValidatorTest, DISABLED_AssignToScalar_Fail) {
// 1 = my_var;
ast::type::I32Type i32;

291
src/writer/hlsl/generator_impl.cc
View File

@ -16,7 +16,6 @@
#include <sstream>
#include "spirv/unified1/GLSL.std.450.h"
#include "src/ast/array_accessor_expression.h"
#include "src/ast/as_expression.h"
#include "src/ast/assignment_statement.h"
@ -31,7 +30,6 @@
#include "src/ast/float_literal.h"
#include "src/ast/identifier_expression.h"
#include "src/ast/if_statement.h"
#include "src/ast/intrinsic.h"
#include "src/ast/loop_statement.h"
#include "src/ast/member_accessor_expression.h"
#include "src/ast/return_statement.h"
@ -469,44 +467,46 @@ bool GeneratorImpl::EmitBreak(std::ostream& out, ast::BreakStatement*) {
return true;
}
std::string GeneratorImpl::generate_intrinsic_name(const std::string& name) {
if (name == "any") {
std::string GeneratorImpl::generate_intrinsic_name(ast::Intrinsic intrinsic) {
if (intrinsic == ast::Intrinsic::kAny) {
return "any";
}
if (name == "all") {
if (intrinsic == ast::Intrinsic::kAll) {
return "all";
}
if (name == "dot") {
if (intrinsic == ast::Intrinsic::kDot) {
return "dot";
}
if (name == "is_finite") {
if (intrinsic == ast::Intrinsic::kIsFinite) {
return "isfinite";
}
if (name == "is_inf") {
if (intrinsic == ast::Intrinsic::kIsInf) {
return "isinf";
}
if (name == "is_nan") {
if (intrinsic == ast::Intrinsic::kIsNan) {
return "isnan";
}
if (name == "dpdy") {
if (intrinsic == ast::Intrinsic::kDpdy) {
return "ddy";
}
if (name == "dpdy_fine") {
if (intrinsic == ast::Intrinsic::kDpdyFine) {
return "ddy_fine";
}
if (name == "dpdy_coarse") {
if (intrinsic == ast::Intrinsic::kDpdyCoarse) {
return "ddy_coarse";
}
if (name == "dpdx") {
if (intrinsic == ast::Intrinsic::kDpdx) {
return "ddx";
}
if (name == "dpdx_fine") {
if (intrinsic == ast::Intrinsic::kDpdxFine) {
return "ddx_fine";
}
if (name == "dpdx_coarse") {
if (intrinsic == ast::Intrinsic::kDpdxCoarse) {
return "ddx_coarse";
}
if (name == "fwidth" || name == "fwidth_fine" || name == "fwidth_coarse") {
if (intrinsic == ast::Intrinsic::kFwidth ||
intrinsic == ast::Intrinsic::kFwidthFine ||
intrinsic == ast::Intrinsic::kFwidthCoarse) {
return "fwidth";
}
return "";
@ -521,15 +521,15 @@ bool GeneratorImpl::EmitCall(std::ostream& pre,
}
auto* ident = expr->func()->AsIdentifier();
if (!ident->has_path() && ast::intrinsic::IsIntrinsic(ident->name())) {
if (ident->IsIntrinsic()) {
const auto& params = expr->params();
if (ident->name() == "select") {
if (ident->intrinsic() == ast::Intrinsic::kSelect) {
error_ = "select not supported in HLSL backend yet";
return false;
} else if (ident->name() == "is_normal") {
} else if (ident->intrinsic() == ast::Intrinsic::kIsNormal) {
error_ = "is_normal not supported in HLSL backend yet";
return false;
} else if (ident->name() == "outer_product") {
} else if (ident->intrinsic() == ast::Intrinsic::kOuterProduct) {
error_ = "outer_product not supported yet";
return false;
// TODO(dsinclair): This gets tricky. We need to generate two variables to
@ -582,11 +582,15 @@ bool GeneratorImpl::EmitCall(std::ostream& pre,
// out << ")";
} else {
auto name = generate_intrinsic_name(ident->name());
auto name = generate_intrinsic_name(ident->intrinsic());
if (name.empty()) {
error_ = "unable to determine intrinsic name for intrinsic: " +
ident->name();
return false;
if (ast::intrinsic::IsTextureIntrinsic(ident->intrinsic())) {
error_ = "Textures not implemented yet";
return false;
}
if (!EmitBuiltinName(pre, out, expr)) {
return false;
}
}
make_indent(out);
@ -609,162 +613,39 @@ bool GeneratorImpl::EmitCall(std::ostream& pre,
return true;
}
if (!ident->has_path()) {
auto name = ident->name();
auto it = ep_func_name_remapped_.find(current_ep_name_ + "_" + name);
if (it != ep_func_name_remapped_.end()) {
name = it->second;
}
auto name = ident->name();
auto it = ep_func_name_remapped_.find(current_ep_name_ + "_" + name);
if (it != ep_func_name_remapped_.end()) {
name = it->second;
}
auto* func = module_->FindFunctionByName(ident->name());
if (func == nullptr) {
error_ = "Unable to find function: " + name;
return false;
}
auto* func = module_->FindFunctionByName(ident->name());
if (func == nullptr) {
error_ = "Unable to find function: " + name;
return false;
}
out << name << "(";
out << name << "(";
bool first = true;
if (has_referenced_in_var_needing_struct(func)) {
auto var_name = current_ep_var_name(VarType::kIn);
if (!var_name.empty()) {
out << var_name;
first = false;
}
bool first = true;
if (has_referenced_in_var_needing_struct(func)) {
auto var_name = current_ep_var_name(VarType::kIn);
if (!var_name.empty()) {
out << var_name;
first = false;
}
if (has_referenced_out_var_needing_struct(func)) {
auto var_name = current_ep_var_name(VarType::kOut);
if (!var_name.empty()) {
if (!first) {
out << ", ";
}
first = false;
out << var_name;
}
}
const auto& params = expr->params();
for (const auto& param : params) {
}
if (has_referenced_out_var_needing_struct(func)) {
auto var_name = current_ep_var_name(VarType::kOut);
if (!var_name.empty()) {
if (!first) {
out << ", ";
}
first = false;
if (!EmitExpression(pre, out, param.get())) {
return false;
}
out << var_name;
}
out << ")";
} else {
return EmitImportFunction(pre, out, expr);
}
return true;
}
bool GeneratorImpl::EmitImportFunction(std::ostream& pre,
std::ostream& out,
ast::CallExpression* expr) {
auto* ident = expr->func()->AsIdentifier();
auto* imp = module_->FindImportByName(ident->path());
if (imp == nullptr) {
error_ = "unable to find import for " + ident->path();
return 0;
}
auto id = imp->GetIdForMethod(ident->name());
if (id == 0) {
error_ = "unable to lookup: " + ident->name() + " in " + ident->path();
}
switch (id) {
case GLSLstd450Acos:
case GLSLstd450Asin:
case GLSLstd450Atan:
case GLSLstd450Atan2:
case GLSLstd450Ceil:
case GLSLstd450Cos:
case GLSLstd450Cosh:
case GLSLstd450Cross:
case GLSLstd450Degrees:
case GLSLstd450Determinant:
case GLSLstd450Distance:
case GLSLstd450Exp:
case GLSLstd450Exp2:
case GLSLstd450FaceForward:
case GLSLstd450Floor:
case GLSLstd450Fma:
case GLSLstd450Length:
case GLSLstd450Log:
case GLSLstd450Log2:
case GLSLstd450Normalize:
case GLSLstd450Pow:
case GLSLstd450Radians:
case GLSLstd450Reflect:
case GLSLstd450Round:
case GLSLstd450Sin:
case GLSLstd450Sinh:
case GLSLstd450SmoothStep:
case GLSLstd450Sqrt:
case GLSLstd450Step:
case GLSLstd450Tan:
case GLSLstd450Tanh:
case GLSLstd450Trunc:
out << ident->name();
break;
case GLSLstd450Fract:
out << "frac";
break;
case GLSLstd450InterpolateAtCentroid:
out << "EvaluateAttributeAtCentroid";
break;
case GLSLstd450InverseSqrt:
out << "rsqrt";
break;
case GLSLstd450FMix:
out << "mix";
break;
case GLSLstd450SSign:
case GLSLstd450FSign:
out << "sign";
break;
case GLSLstd450FAbs:
case GLSLstd450SAbs:
out << "abs";
break;
case GLSLstd450FMax:
case GLSLstd450NMax:
case GLSLstd450SMax:
case GLSLstd450UMax:
out << "max";
break;
case GLSLstd450FMin:
case GLSLstd450NMin:
case GLSLstd450SMin:
case GLSLstd450UMin:
out << "min";
break;
case GLSLstd450FClamp:
case GLSLstd450SClamp:
case GLSLstd450NClamp:
case GLSLstd450UClamp:
out << "clamp";
break;
// TODO(dsinclair): Determine mappings for the following
case GLSLstd450Atanh:
case GLSLstd450Asinh:
case GLSLstd450Acosh:
case GLSLstd450FindILsb:
case GLSLstd450FindUMsb:
case GLSLstd450FindSMsb:
case GLSLstd450MatrixInverse:
case GLSLstd450RoundEven:
error_ = "Unknown import method: " + ident->name();
return false;
}
out << "(";
bool first = true;
const auto& params = expr->params();
for (const auto& param : params) {
if (!first) {
@ -776,11 +657,73 @@ bool GeneratorImpl::EmitImportFunction(std::ostream& pre,
return false;
}
}
out << ")";
return true;
}
bool GeneratorImpl::EmitBuiltinName(std::ostream&,
std::ostream& out,
ast::CallExpression* expr) {
auto* ident = expr->func()->AsIdentifier();
switch (ident->intrinsic()) {
case ast::Intrinsic::kAcos:
case ast::Intrinsic::kAsin:
case ast::Intrinsic::kAtan:
case ast::Intrinsic::kAtan2:
case ast::Intrinsic::kCeil:
case ast::Intrinsic::kCos:
case ast::Intrinsic::kCosh:
case ast::Intrinsic::kCross:
case ast::Intrinsic::kDeterminant:
case ast::Intrinsic::kDistance:
case ast::Intrinsic::kExp:
case ast::Intrinsic::kExp2:
case ast::Intrinsic::kFloor:
case ast::Intrinsic::kFma:
case ast::Intrinsic::kLength:
case ast::Intrinsic::kLog:
case ast::Intrinsic::kLog2:
case ast::Intrinsic::kNormalize:
case ast::Intrinsic::kPow:
case ast::Intrinsic::kReflect:
case ast::Intrinsic::kRound:
case ast::Intrinsic::kSin:
case ast::Intrinsic::kSinh:
case ast::Intrinsic::kSqrt:
case ast::Intrinsic::kStep:
case ast::Intrinsic::kTan:
case ast::Intrinsic::kTanh:
case ast::Intrinsic::kTrunc:
case ast::Intrinsic::kMix:
case ast::Intrinsic::kSign:
case ast::Intrinsic::kAbs:
case ast::Intrinsic::kMax:
case ast::Intrinsic::kMin:
case ast::Intrinsic::kClamp:
out << ident->name();
break;
case ast::Intrinsic::kFaceForward:
out << "faceforward";
break;
case ast::Intrinsic::kFract:
out << "frac";
break;
case ast::Intrinsic::kInverseSqrt:
out << "rsqrt";
break;
case ast::Intrinsic::kSmoothStep:
out << "smoothstep";
break;
default:
error_ = "Unknown builtin method: " + ident->name();
return false;
}
return true;
}
bool GeneratorImpl::EmitCast(std::ostream& pre,
std::ostream& out,
ast::CastExpression* expr) {
@ -957,12 +900,6 @@ bool GeneratorImpl::EmitIdentifier(std::ostream&,
std::ostream& out,
ast::IdentifierExpression* expr) {
auto* ident = expr->AsIdentifier();
if (ident->has_path()) {
// TODO(dsinclair): Handle identifier with path
error_ = "Identifier paths not handled yet.";
return false;
}
ast::Variable* var = nullptr;
if (global_variables_.get(ident->name(), &var)) {
if (global_is_in_struct(var)) {
@ -1832,10 +1769,6 @@ bool GeneratorImpl::is_storage_buffer_access(
// Check if this is a storage buffer variable
if (structure->IsIdentifier()) {
auto* ident = expr->structure()->AsIdentifier();
if (ident->has_path()) {
return false;
}
ast::Variable* var = nullptr;
if (!global_variables_.get(ident->name(), &var)) {
return false;

15
src/writer/hlsl/generator_impl.h
View File

@ -15,6 +15,7 @@
#ifndef SRC_WRITER_HLSL_GENERATOR_IMPL_H_
#define SRC_WRITER_HLSL_GENERATOR_IMPL_H_
#include "src/ast/intrinsic.h"
#include "src/ast/literal.h"
#include "src/ast/module.h"
#include "src/ast/scalar_constructor_expression.h"
@ -206,14 +207,14 @@ class GeneratorImpl {
/// @param stmt the statement to emit
/// @returns true if the statement was successfully emitted
bool EmitIf(std::ostream& out, ast::IfStatement* stmt);
/// Handles genreating an import expression
/// Handles generating a builtin method name
/// @param pre the preamble for the expression stream
/// @param out the output of the expression stream
/// @param expr the expression
/// @returns true if the expression was successfully emitted.
bool EmitImportFunction(std::ostream& pre,
std::ostream& out,
ast::CallExpression* expr);
/// @returns true if the name was successfully emitted.
bool EmitBuiltinName(std::ostream& pre,
std::ostream& out,
ast::CallExpression* expr);
/// Handles a literal
/// @param out the output stream
/// @param lit the literal to emit
@ -327,9 +328,9 @@ class GeneratorImpl {
/// @returns the name
std::string generate_name(const std::string& prefix);
/// Generates an intrinsic name from the given name
/// @param name the name to convert to an intrinsic
/// @param intrinsic the intrinsic to convert to a name
/// @returns the intrinsic name or blank on error
std::string generate_intrinsic_name(const std::string& name);
std::string generate_intrinsic_name(ast::Intrinsic intrinsic);
/// Converts a builtin to an attribute name
/// @param builtin the builtin to convert
/// @returns the string name of the builtin or blank on error

16
src/writer/hlsl/generator_impl_identifier_test.cc
View File

@ -23,13 +23,7 @@ namespace {
using HlslGeneratorImplTest_Identifier = TestHelper;
TEST_F(HlslGeneratorImplTest_Identifier, DISABLED_EmitExpression_Identifier) {
ast::IdentifierExpression i(std::vector<std::string>{"std", "glsl"});
ASSERT_TRUE(gen().EmitExpression(pre(), out(), &i)) << gen().error();
EXPECT_EQ(result(), "std::glsl");
}
TEST_F(HlslGeneratorImplTest_Identifier, EmitIdentifierExpression_Single) {
TEST_F(HlslGeneratorImplTest_Identifier, EmitIdentifierExpression) {
ast::IdentifierExpression i("foo");
ASSERT_TRUE(gen().EmitExpression(pre(), out(), &i)) << gen().error();
EXPECT_EQ(result(), "foo");
@ -42,14 +36,6 @@ TEST_F(HlslGeneratorImplTest_Identifier,
EXPECT_EQ(result(), "virtual_tint_0");
}
// TODO(dsinclair): Handle import names
TEST_F(HlslGeneratorImplTest_Identifier,
DISABLED_EmitIdentifierExpression_MultipleNames) {
ast::IdentifierExpression i({"std", "glsl", "init"});
ASSERT_TRUE(gen().EmitExpression(pre(), out(), &i)) << gen().error();
EXPECT_EQ(result(), "std::glsl::init");
}
} // namespace
} // namespace hlsl
} // namespace writer

173
src/writer/hlsl/generator_impl_import_test.cc
View File

@ -58,59 +58,40 @@ TEST_P(HlslImportData_SingleParamTest, FloatScalar) {
params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::FloatLiteral>(&f32, 1.f)));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", param.name}),
std::move(params));
mod()->AddImport(std::make_unique<ast::Import>("GLSL.std.450", "std"));
auto ident = std::make_unique<ast::IdentifierExpression>(param.name);
ast::CallExpression expr(std::move(ident), std::move(params));
ASSERT_TRUE(td().DetermineResultType(&expr)) << td().error();
ASSERT_TRUE(gen().EmitImportFunction(pre(), out(), &expr)) << gen().error();
ASSERT_TRUE(gen().EmitCall(pre(), out(), &expr)) << gen().error();
EXPECT_EQ(result(), std::string(param.hlsl_name) + "(1.00000000f)");
}
INSTANTIATE_TEST_SUITE_P(
HlslGeneratorImplTest_Import,
HlslImportData_SingleParamTest,
testing::Values(HlslImportData{"acos", "acos"},
HlslImportData{"asin", "asin"},
HlslImportData{"atan", "atan"},
HlslImportData{"cos", "cos"},
HlslImportData{"cosh", "cosh"},
HlslImportData{"ceil", "ceil"},
HlslImportData{"degrees", "degrees"},
HlslImportData{"exp", "exp"},
HlslImportData{"exp2", "exp2"},
HlslImportData{"fabs", "abs"},
HlslImportData{"floor", "floor"},
HlslImportData{"fract", "frac"},
HlslImportData{"interpolateatcentroid",
"EvaluateAttributeAtCentroid"},
HlslImportData{"inversesqrt", "rsqrt"},
HlslImportData{"length", "length"},
HlslImportData{"log", "log"},
HlslImportData{"log2", "log2"},
HlslImportData{"normalize", "normalize"},
HlslImportData{"radians", "radians"},
HlslImportData{"round", "round"},
HlslImportData{"fsign", "sign"},
HlslImportData{"sin", "sin"},
HlslImportData{"sinh", "sinh"},
HlslImportData{"sqrt", "sqrt"},
HlslImportData{"tan", "tan"},
HlslImportData{"tanh", "tanh"},
HlslImportData{"trunc", "trunc"}));
TEST_F(HlslGeneratorImplTest_Import, DISABLED_HlslImportData_Acosh) {
FAIL();
}
TEST_F(HlslGeneratorImplTest_Import, DISABLED_HlslImportData_ASinh) {
FAIL();
}
TEST_F(HlslGeneratorImplTest_Import, DISABLED_HlslImportData_ATanh) {
FAIL();
}
INSTANTIATE_TEST_SUITE_P(HlslGeneratorImplTest_Import,
HlslImportData_SingleParamTest,
testing::Values(HlslImportData{"abs", "abs"},
HlslImportData{"acos", "acos"},
HlslImportData{"asin", "asin"},
HlslImportData{"atan", "atan"},
HlslImportData{"cos", "cos"},
HlslImportData{"cosh", "cosh"},
HlslImportData{"ceil", "ceil"},
HlslImportData{"exp", "exp"},
HlslImportData{"exp2", "exp2"},
HlslImportData{"floor", "floor"},
HlslImportData{"fract", "frac"},
HlslImportData{"inverseSqrt", "rsqrt"},
HlslImportData{"length", "length"},
HlslImportData{"log", "log"},
HlslImportData{"log2", "log2"},
HlslImportData{"normalize",
"normalize"},
HlslImportData{"round", "round"},
HlslImportData{"sign", "sign"},
HlslImportData{"sin", "sin"},
HlslImportData{"sinh", "sinh"},
HlslImportData{"sqrt", "sqrt"},
HlslImportData{"tan", "tan"},
HlslImportData{"tanh", "tanh"},
HlslImportData{"trunc", "trunc"}));
using HlslImportData_SingleIntParamTest =
TestHelperBase<testing::TestWithParam<HlslImportData>>;
@ -123,20 +104,17 @@ TEST_P(HlslImportData_SingleIntParamTest, IntScalar) {
params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::SintLiteral>(&i32, 1)));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", param.name}),
std::move(params));
mod()->AddImport(std::make_unique<ast::Import>("GLSL.std.450", "std"));
ast::CallExpression expr(
std::make_unique<ast::IdentifierExpression>(param.name),
std::move(params));
ASSERT_TRUE(td().DetermineResultType(&expr)) << td().error();
ASSERT_TRUE(gen().EmitImportFunction(pre(), out(), &expr)) << gen().error();
ASSERT_TRUE(gen().EmitCall(pre(), out(), &expr)) << gen().error();
EXPECT_EQ(result(), std::string(param.hlsl_name) + "(1)");
}
INSTANTIATE_TEST_SUITE_P(HlslGeneratorImplTest_Import,
HlslImportData_SingleIntParamTest,
testing::Values(HlslImportData{"sabs", "abs"},
HlslImportData{"ssign", "sign"}));
testing::Values(HlslImportData{"abs", "abs"}));
using HlslImportData_DualParamTest =
TestHelperBase<testing::TestWithParam<HlslImportData>>;
@ -151,14 +129,12 @@ TEST_P(HlslImportData_DualParamTest, FloatScalar) {
params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::FloatLiteral>(&f32, 2.f)));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", param.name}),
std::move(params));
mod()->AddImport(std::make_unique<ast::Import>("GLSL.std.450", "std"));
ast::CallExpression expr(
std::make_unique<ast::IdentifierExpression>(param.name),
std::move(params));
ASSERT_TRUE(td().DetermineResultType(&expr)) << td().error();
ASSERT_TRUE(gen().EmitImportFunction(pre(), out(), &expr)) << gen().error();
ASSERT_TRUE(gen().EmitCall(pre(), out(), &expr)) << gen().error();
EXPECT_EQ(result(),
std::string(param.hlsl_name) + "(1.00000000f, 2.00000000f)");
}
@ -166,10 +142,8 @@ INSTANTIATE_TEST_SUITE_P(HlslGeneratorImplTest_Import,
HlslImportData_DualParamTest,
testing::Values(HlslImportData{"atan2", "atan2"},
HlslImportData{"distance", "distance"},
HlslImportData{"fmax", "max"},
HlslImportData{"fmin", "min"},
HlslImportData{"nmax", "max"},
HlslImportData{"nmin", "min"},
HlslImportData{"max", "max"},
HlslImportData{"min", "min"},
HlslImportData{"pow", "pow"},
HlslImportData{"reflect", "reflect"},
HlslImportData{"step", "step"}));
@ -203,14 +177,12 @@ TEST_P(HlslImportData_DualParam_VectorTest, FloatVector) {
params.push_back(std::make_unique<ast::TypeConstructorExpression>(
&vec, std::move(type_params)));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", param.name}),
std::move(params));
mod()->AddImport(std::make_unique<ast::Import>("GLSL.std.450", "std"));
ast::CallExpression expr(
std::make_unique<ast::IdentifierExpression>(param.name),
std::move(params));
ASSERT_TRUE(td().DetermineResultType(&expr)) << td().error();
ASSERT_TRUE(gen().EmitImportFunction(pre(), out(), &expr)) << gen().error();
ASSERT_TRUE(gen().EmitCall(pre(), out(), &expr)) << gen().error();
EXPECT_EQ(result(),
std::string(param.hlsl_name) +
"(vector<float, 3>(1.00000000f, 2.00000000f, 3.00000000f), "
@ -233,22 +205,18 @@ TEST_P(HlslImportData_DualParam_Int_Test, IntScalar) {
params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::SintLiteral>(&i32, 2)));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", param.name}),
std::move(params));
mod()->AddImport(std::make_unique<ast::Import>("GLSL.std.450", "std"));
ast::CallExpression expr(
std::make_unique<ast::IdentifierExpression>(param.name),
std::move(params));
ASSERT_TRUE(td().DetermineResultType(&expr)) << td().error();
ASSERT_TRUE(gen().EmitImportFunction(pre(), out(), &expr)) << gen().error();
ASSERT_TRUE(gen().EmitCall(pre(), out(), &expr)) << gen().error();
EXPECT_EQ(result(), std::string(param.hlsl_name) + "(1, 2)");
}
INSTANTIATE_TEST_SUITE_P(HlslGeneratorImplTest_Import,
HlslImportData_DualParam_Int_Test,
testing::Values(HlslImportData{"smax", "max"},
HlslImportData{"smin", "min"},
HlslImportData{"umax", "max"},
HlslImportData{"umin", "min"}));
testing::Values(HlslImportData{"max", "max"},
HlslImportData{"min", "min"}));
using HlslImportData_TripleParamTest =
TestHelperBase<testing::TestWithParam<HlslImportData>>;
@ -265,25 +233,22 @@ TEST_P(HlslImportData_TripleParamTest, FloatScalar) {
params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::FloatLiteral>(&f32, 3.f)));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", param.name}),
std::move(params));
mod()->AddImport(std::make_unique<ast::Import>("GLSL.std.450", "std"));
ast::CallExpression expr(
std::make_unique<ast::IdentifierExpression>(param.name),
std::move(params));
ASSERT_TRUE(td().DetermineResultType(&expr)) << td().error();
ASSERT_TRUE(gen().EmitImportFunction(pre(), out(), &expr)) << gen().error();
ASSERT_TRUE(gen().EmitCall(pre(), out(), &expr)) << gen().error();
EXPECT_EQ(result(), std::string(param.hlsl_name) +
"(1.00000000f, 2.00000000f, 3.00000000f)");
}
INSTANTIATE_TEST_SUITE_P(
HlslGeneratorImplTest_Import,
HlslImportData_TripleParamTest,
testing::Values(HlslImportData{"faceforward", "faceforward"},
testing::Values(HlslImportData{"faceForward", "faceforward"},
HlslImportData{"fma", "fma"},
HlslImportData{"fclamp", "clamp"},
HlslImportData{"nclamp", "clamp"},
HlslImportData{"smoothstep", "smoothstep"}));
HlslImportData{"clamp", "clamp"},
HlslImportData{"smoothStep", "smoothstep"}));
TEST_F(HlslGeneratorImplTest_Import, DISABLED_HlslImportData_FMix) {
FAIL();
@ -304,20 +269,17 @@ TEST_P(HlslImportData_TripleParam_Int_Test, IntScalar) {
params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::SintLiteral>(&i32, 3)));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", param.name}),
std::move(params));
mod()->AddImport(std::make_unique<ast::Import>("GLSL.std.450", "std"));
ast::CallExpression expr(
std::make_unique<ast::IdentifierExpression>(param.name),
std::move(params));
ASSERT_TRUE(td().DetermineResultType(&expr)) << td().error();
ASSERT_TRUE(gen().EmitImportFunction(pre(), out(), &expr)) << gen().error();
ASSERT_TRUE(gen().EmitCall(pre(), out(), &expr)) << gen().error();
EXPECT_EQ(result(), std::string(param.hlsl_name) + "(1, 2, 3)");
}
INSTANTIATE_TEST_SUITE_P(HlslGeneratorImplTest_Import,
HlslImportData_TripleParam_Int_Test,
testing::Values(HlslImportData{"sclamp", "clamp"},
HlslImportData{"uclamp", "clamp"}));
testing::Values(HlslImportData{"clamp", "clamp"}));
TEST_F(HlslGeneratorImplTest_Import, HlslImportData_Determinant) {
ast::type::F32Type f32;
@ -329,17 +291,16 @@ TEST_F(HlslGeneratorImplTest_Import, HlslImportData_Determinant) {
ast::ExpressionList params;
params.push_back(std::make_unique<ast::IdentifierExpression>("var"));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", "determinant"}),
std::move(params));
ast::CallExpression expr(
std::make_unique<ast::IdentifierExpression>("determinant"),
std::move(params));
mod()->AddGlobalVariable(std::move(var));
mod()->AddImport(std::make_unique<ast::Import>("GLSL.std.450", "std"));
// Register the global
ASSERT_TRUE(td().Determine()) << td().error();
ASSERT_TRUE(td().DetermineResultType(&expr)) << td().error();
ASSERT_TRUE(gen().EmitImportFunction(pre(), out(), &expr)) << gen().error();
ASSERT_TRUE(gen().EmitCall(pre(), out(), &expr)) << gen().error();
EXPECT_EQ(result(), std::string("determinant(var)"));
}

40
src/writer/hlsl/generator_impl_intrinsic_test.cc
View File

@ -29,36 +29,36 @@ namespace {
using HlslGeneratorImplTest_Intrinsic = TestHelper;
struct IntrinsicData {
const char* name;
ast::Intrinsic intrinsic;
const char* hlsl_name;
};
inline std::ostream& operator<<(std::ostream& out, IntrinsicData data) {
out << data.name;
out << data.hlsl_name;
return out;
}
using HlslIntrinsicTest = TestHelperBase<testing::TestWithParam<IntrinsicData>>;
TEST_P(HlslIntrinsicTest, Emit) {
auto param = GetParam();
EXPECT_EQ(gen().generate_intrinsic_name(param.name), param.hlsl_name);
EXPECT_EQ(gen().generate_intrinsic_name(param.intrinsic), param.hlsl_name);
}
INSTANTIATE_TEST_SUITE_P(
HlslGeneratorImplTest_Intrinsic,
HlslIntrinsicTest,
testing::Values(IntrinsicData{"any", "any"},
IntrinsicData{"all", "all"},
IntrinsicData{"dot", "dot"},
IntrinsicData{"dpdx", "ddx"},
IntrinsicData{"dpdx_coarse", "ddx_coarse"},
IntrinsicData{"dpdx_fine", "ddx_fine"},
IntrinsicData{"dpdy", "ddy"},
IntrinsicData{"dpdy_coarse", "ddy_coarse"},
IntrinsicData{"dpdy_fine", "ddy_fine"},
IntrinsicData{"fwidth", "fwidth"},
IntrinsicData{"fwidth_coarse", "fwidth"},
IntrinsicData{"fwidth_fine", "fwidth"},
IntrinsicData{"is_finite", "isfinite"},
IntrinsicData{"is_inf", "isinf"},
IntrinsicData{"is_nan", "isnan"}));
testing::Values(IntrinsicData{ast::Intrinsic::kAny, "any"},
IntrinsicData{ast::Intrinsic::kAll, "all"},
IntrinsicData{ast::Intrinsic::kDot, "dot"},
IntrinsicData{ast::Intrinsic::kDpdx, "ddx"},
IntrinsicData{ast::Intrinsic::kDpdxCoarse, "ddx_coarse"},
IntrinsicData{ast::Intrinsic::kDpdxFine, "ddx_fine"},
IntrinsicData{ast::Intrinsic::kDpdy, "ddy"},
IntrinsicData{ast::Intrinsic::kDpdyCoarse, "ddy_coarse"},
IntrinsicData{ast::Intrinsic::kDpdyFine, "ddy_fine"},
IntrinsicData{ast::Intrinsic::kFwidth, "fwidth"},
IntrinsicData{ast::Intrinsic::kFwidthCoarse, "fwidth"},
IntrinsicData{ast::Intrinsic::kFwidthFine, "fwidth"},
IntrinsicData{ast::Intrinsic::kIsFinite, "isfinite"},
IntrinsicData{ast::Intrinsic::kIsInf, "isinf"},
IntrinsicData{ast::Intrinsic::kIsNan, "isnan"}));
TEST_F(HlslGeneratorImplTest_Intrinsic, DISABLED_Intrinsic_IsNormal) {
FAIL();
@ -101,7 +101,7 @@ TEST_F(HlslGeneratorImplTest_Intrinsic, DISABLED_Intrinsic_OuterProduct) {
}
TEST_F(HlslGeneratorImplTest_Intrinsic, Intrinsic_Bad_Name) {
EXPECT_EQ(gen().generate_intrinsic_name("unknown name"), "");
EXPECT_EQ(gen().generate_intrinsic_name(ast::Intrinsic::kNone), "");
}
TEST_F(HlslGeneratorImplTest_Intrinsic, Intrinsic_Call) {
@ -112,6 +112,8 @@ TEST_F(HlslGeneratorImplTest_Intrinsic, Intrinsic_Call) {
ast::CallExpression call(std::make_unique<ast::IdentifierExpression>("dot"),
std::move(params));
ASSERT_TRUE(td().DetermineResultType(&call)) << td().error();
gen().increment_indent();
ASSERT_TRUE(gen().EmitExpression(pre(), out(), &call)) << gen().error();
EXPECT_EQ(result(), " dot(param1, param2)");

370
src/writer/msl/generator_impl.cc
View File

@ -14,7 +14,6 @@
#include "src/writer/msl/generator_impl.h"
#include "spirv/unified1/GLSL.std.450.h"
#include "src/ast/array_accessor_expression.h"
#include "src/ast/as_expression.h"
#include "src/ast/assignment_statement.h"
@ -33,7 +32,6 @@
#include "src/ast/function.h"
#include "src/ast/identifier_expression.h"
#include "src/ast/if_statement.h"
#include "src/ast/intrinsic.h"
#include "src/ast/location_decoration.h"
#include "src/ast/loop_statement.h"
#include "src/ast/member_accessor_expression.h"
@ -409,38 +407,44 @@ std::string GeneratorImpl::current_ep_var_name(VarType type) {
return name;
}
std::string GeneratorImpl::generate_intrinsic_name(const std::string& name) {
if (name == "any") {
std::string GeneratorImpl::generate_intrinsic_name(ast::Intrinsic intrinsic) {
if (intrinsic == ast::Intrinsic::kAny) {
return "any";
}
if (name == "all") {
if (intrinsic == ast::Intrinsic::kAll) {
return "all";
}
if (name == "dot") {
if (intrinsic == ast::Intrinsic::kDot) {
return "dot";
}
if (name == "is_finite") {
if (intrinsic == ast::Intrinsic::kIsFinite) {
return "isfinite";
}
if (name == "is_inf") {
if (intrinsic == ast::Intrinsic::kIsInf) {
return "isinf";
}
if (name == "is_nan") {
if (intrinsic == ast::Intrinsic::kIsNan) {
return "isnan";
}
if (name == "is_normal") {
if (intrinsic == ast::Intrinsic::kIsNormal) {
return "isnormal";
}
if (name == "select") {
if (intrinsic == ast::Intrinsic::kSelect) {
return "select";
}
if (name == "dpdy" || name == "dpdy_fine" || name == "dpdy_coarse") {
if (intrinsic == ast::Intrinsic::kDpdy ||
intrinsic == ast::Intrinsic::kDpdyFine ||
intrinsic == ast::Intrinsic::kDpdyCoarse) {
return "dfdy";
}
if (name == "dpdx" || name == "dpdx_fine" || name == "dpdx_coarse") {
if (intrinsic == ast::Intrinsic::kDpdx ||
intrinsic == ast::Intrinsic::kDpdxFine ||
intrinsic == ast::Intrinsic::kDpdxCoarse) {
return "dfdx";
}
if (name == "fwidth" || name == "fwidth_fine" || name == "fwidth_coarse") {
if (intrinsic == ast::Intrinsic::kFwidth ||
intrinsic == ast::Intrinsic::kFwidthFine ||
intrinsic == ast::Intrinsic::kFwidthCoarse) {
return "fwidth";
}
return "";
@ -453,9 +457,9 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) {
}
auto* ident = expr->func()->AsIdentifier();
if (!ident->has_path() && ast::intrinsic::IsIntrinsic(ident->name())) {
if (ident->IsIntrinsic()) {
const auto& params = expr->params();
if (ident->name() == "outer_product") {
if (ident->intrinsic() == ast::Intrinsic::kOuterProduct) {
error_ = "outer_product not supported yet";
return false;
// TODO(dsinclair): This gets tricky. We need to generate two variables to
@ -508,11 +512,15 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) {
// out_ << ")";
} else {
auto name = generate_intrinsic_name(ident->name());
auto name = generate_intrinsic_name(ident->intrinsic());
if (name.empty()) {
error_ = "unable to determine intrinsic name for intrinsic: " +
ident->name();
return false;
if (ast::intrinsic::IsTextureIntrinsic(ident->intrinsic())) {
error_ = "Textures not implemented yet";
return false;
}
if (!EmitBuiltinName(ident)) {
return false;
}
}
make_indent();
@ -535,191 +543,69 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) {
return true;
}
if (!ident->has_path()) {
auto name = ident->name();
auto it = ep_func_name_remapped_.find(current_ep_name_ + "_" + name);
if (it != ep_func_name_remapped_.end()) {
name = it->second;
}
auto* func = module_->FindFunctionByName(ident->name());
if (func == nullptr) {
error_ = "Unable to find function: " + name;
return false;
}
out_ << name << "(";
bool first = true;
if (has_referenced_in_var_needing_struct(func)) {
auto var_name = current_ep_var_name(VarType::kIn);
if (!var_name.empty()) {
out_ << var_name;
first = false;
}
}
if (has_referenced_out_var_needing_struct(func)) {
auto var_name = current_ep_var_name(VarType::kOut);
if (!var_name.empty()) {
if (!first) {
out_ << ", ";
}
first = false;
out_ << var_name;
}
}
for (const auto& data : func->referenced_builtin_variables()) {
auto* var = data.first;
if (var->storage_class() != ast::StorageClass::kInput) {
continue;
}
if (!first) {
out_ << ", ";
}
first = false;
out_ << var->name();
}
for (const auto& data : func->referenced_uniform_variables()) {
auto* var = data.first;
if (!first) {
out_ << ", ";
}
first = false;
out_ << var->name();
}
for (const auto& data : func->referenced_storagebuffer_variables()) {
auto* var = data.first;
if (!first) {
out_ << ", ";
}
first = false;
out_ << var->name();
}
const auto& params = expr->params();
for (const auto& param : params) {
if (!first) {
out_ << ", ";
}
first = false;
if (!EmitExpression(param.get())) {
return false;
}
}
out_ << ")";
} else {
return EmitImportFunction(expr);
}
return true;
}
bool GeneratorImpl::EmitImportFunction(ast::CallExpression* expr) {
auto* ident = expr->func()->AsIdentifier();
auto* imp = module_->FindImportByName(ident->path());
if (imp == nullptr) {
error_ = "unable to find import for " + ident->path();
return 0;
}
auto id = imp->GetIdForMethod(ident->name());
if (id == 0) {
error_ = "unable to lookup: " + ident->name() + " in " + ident->path();
auto name = ident->name();
auto it = ep_func_name_remapped_.find(current_ep_name_ + "_" + name);
if (it != ep_func_name_remapped_.end()) {
name = it->second;
}
out_ << "metal::";
switch (id) {
case GLSLstd450Acos:
case GLSLstd450Acosh:
case GLSLstd450Asin:
case GLSLstd450Asinh:
case GLSLstd450Atan:
case GLSLstd450Atan2:
case GLSLstd450Atanh:
case GLSLstd450Ceil:
case GLSLstd450Cos:
case GLSLstd450Cosh:
case GLSLstd450Cross:
case GLSLstd450Determinant:
case GLSLstd450Distance:
case GLSLstd450Exp:
case GLSLstd450Exp2:
case GLSLstd450FAbs:
case GLSLstd450FaceForward:
case GLSLstd450Floor:
case GLSLstd450Fma:
case GLSLstd450Fract:
case GLSLstd450Length:
case GLSLstd450Log:
case GLSLstd450Log2:
case GLSLstd450Normalize:
case GLSLstd450Pow:
case GLSLstd450Reflect:
case GLSLstd450Round:
case GLSLstd450Sin:
case GLSLstd450Sinh:
case GLSLstd450SmoothStep:
case GLSLstd450Sqrt:
case GLSLstd450Step:
case GLSLstd450Tan:
case GLSLstd450Tanh:
case GLSLstd450Trunc:
out_ << ident->name();
break;
case GLSLstd450InverseSqrt:
out_ << "rsqrt";
break;
case GLSLstd450FMax:
case GLSLstd450NMax:
out_ << "fmax";
break;
case GLSLstd450FMin:
case GLSLstd450NMin:
out_ << "fmin";
break;
case GLSLstd450FMix:
out_ << "mix";
break;
case GLSLstd450FSign:
out_ << "sign";
break;
case GLSLstd450SAbs:
out_ << "abs";
break;
case GLSLstd450SMax:
case GLSLstd450UMax:
out_ << "max";
break;
case GLSLstd450SMin:
case GLSLstd450UMin:
out_ << "min";
break;
case GLSLstd450FClamp:
case GLSLstd450SClamp:
case GLSLstd450NClamp:
case GLSLstd450UClamp:
out_ << "clamp";
break;
// TODO(dsinclair): Determine mappings for the following
case GLSLstd450Degrees:
case GLSLstd450FindILsb:
case GLSLstd450FindUMsb:
case GLSLstd450FindSMsb:
case GLSLstd450InterpolateAtCentroid:
case GLSLstd450MatrixInverse:
case GLSLstd450Radians:
case GLSLstd450RoundEven:
case GLSLstd450SSign:
error_ = "Unknown import method: " + ident->name();
return false;
auto* func = module_->FindFunctionByName(ident->name());
if (func == nullptr) {
error_ = "Unable to find function: " + name;
return false;
}
out_ << "(";
out_ << name << "(";
bool first = true;
if (has_referenced_in_var_needing_struct(func)) {
auto var_name = current_ep_var_name(VarType::kIn);
if (!var_name.empty()) {
out_ << var_name;
first = false;
}
}
if (has_referenced_out_var_needing_struct(func)) {
auto var_name = current_ep_var_name(VarType::kOut);
if (!var_name.empty()) {
if (!first) {
out_ << ", ";
}
first = false;
out_ << var_name;
}
}
for (const auto& data : func->referenced_builtin_variables()) {
auto* var = data.first;
if (var->storage_class() != ast::StorageClass::kInput) {
continue;
}
if (!first) {
out_ << ", ";
}
first = false;
out_ << var->name();
}
for (const auto& data : func->referenced_uniform_variables()) {
auto* var = data.first;
if (!first) {
out_ << ", ";
}
first = false;
out_ << var->name();
}
for (const auto& data : func->referenced_storagebuffer_variables()) {
auto* var = data.first;
if (!first) {
out_ << ", ";
}
first = false;
out_ << var->name();
}
const auto& params = expr->params();
for (const auto& param : params) {
if (!first) {
@ -731,11 +617,89 @@ bool GeneratorImpl::EmitImportFunction(ast::CallExpression* expr) {
return false;
}
}
out_ << ")";
return true;
}
bool GeneratorImpl::EmitBuiltinName(ast::IdentifierExpression* ident) {
out_ << "metal::";
switch (ident->intrinsic()) {
case ast::Intrinsic::kAcos:
case ast::Intrinsic::kAsin:
case ast::Intrinsic::kAtan:
case ast::Intrinsic::kAtan2:
case ast::Intrinsic::kCeil:
case ast::Intrinsic::kCos:
case ast::Intrinsic::kCosh:
case ast::Intrinsic::kCross:
case ast::Intrinsic::kDeterminant:
case ast::Intrinsic::kDistance:
case ast::Intrinsic::kExp:
case ast::Intrinsic::kExp2:
case ast::Intrinsic::kFloor:
case ast::Intrinsic::kFma:
case ast::Intrinsic::kFract:
case ast::Intrinsic::kLength:
case ast::Intrinsic::kLog:
case ast::Intrinsic::kLog2:
case ast::Intrinsic::kMix:
case ast::Intrinsic::kNormalize:
case ast::Intrinsic::kPow:
case ast::Intrinsic::kReflect:
case ast::Intrinsic::kRound:
case ast::Intrinsic::kSin:
case ast::Intrinsic::kSinh:
case ast::Intrinsic::kSqrt:
case ast::Intrinsic::kStep:
case ast::Intrinsic::kTan:
case ast::Intrinsic::kTanh:
case ast::Intrinsic::kTrunc:
case ast::Intrinsic::kSign:
case ast::Intrinsic::kClamp:
out_ << ident->name();
break;
case ast::Intrinsic::kAbs:
if (ident->result_type()->IsF32()) {
out_ << "fabs";
} else if (ident->result_type()->IsU32() ||
ident->result_type()->IsI32()) {
out_ << "abs";
}
break;
case ast::Intrinsic::kMax:
if (ident->result_type()->IsF32()) {
out_ << "fmax";
} else if (ident->result_type()->IsU32() ||
ident->result_type()->IsI32()) {
out_ << "max";
}
break;
case ast::Intrinsic::kMin:
if (ident->result_type()->IsF32()) {
out_ << "fmin";
} else if (ident->result_type()->IsU32() ||
ident->result_type()->IsI32()) {
out_ << "min";
}
break;
case ast::Intrinsic::kFaceForward:
out_ << "faceforward";
break;
case ast::Intrinsic::kSmoothStep:
out_ << "smoothstep";
break;
case ast::Intrinsic::kInverseSqrt:
out_ << "rsqrt";
break;
default:
error_ = "Unknown import method: " + ident->name();
return false;
}
return true;
}
bool GeneratorImpl::EmitCase(ast::CaseStatement* stmt) {
make_indent();
@ -1402,12 +1366,6 @@ bool GeneratorImpl::global_is_in_struct(ast::Variable* var) const {
bool GeneratorImpl::EmitIdentifier(ast::IdentifierExpression* expr) {
auto* ident = expr->AsIdentifier();
if (ident->has_path()) {
// TODO(dsinclair): Handle identifier with path
error_ = "Identifier paths not handled yet.";
return false;
}
ast::Variable* var = nullptr;
if (global_variables_.get(ident->name(), &var)) {
if (global_is_in_struct(var)) {

13
src/writer/msl/generator_impl.h
View File

@ -19,6 +19,7 @@
#include <string>
#include <unordered_map>
#include "src/ast/intrinsic.h"
#include "src/ast/literal.h"
#include "src/ast/module.h"
#include "src/ast/scalar_constructor_expression.h"
@ -150,10 +151,10 @@ class GeneratorImpl : public TextGenerator {
/// @param stmt the statement to emit
/// @returns true if the statement was successfully emitted
bool EmitIf(ast::IfStatement* stmt);
/// Handles genreating an import expression
/// @param expr the expression
/// @returns true if the expression was successfully emitted.
bool EmitImportFunction(ast::CallExpression* expr);
/// Handles generating a builtin name
/// @param ident the identifier to build the name from
/// @returns true if the name was successfully emitted.
bool EmitBuiltinName(ast::IdentifierExpression* ident);
/// Handles a literal
/// @param lit the literal to emit
/// @returns true if the literal was successfully emitted
@ -230,9 +231,9 @@ class GeneratorImpl : public TextGenerator {
/// @returns the name
std::string generate_name(const std::string& prefix);
/// Generates an intrinsic name from the given name
/// @param name the name to convert to an intrinsic
/// @param intrinsic the intrinsic to convert to an method name
/// @returns the intrinsic name or blank on error
std::string generate_intrinsic_name(const std::string& name);
std::string generate_intrinsic_name(ast::Intrinsic intrinsic);
/// Checks if the global variable is in an input or output struct
/// @param var the variable to check

21
src/writer/msl/generator_impl_identifier_test.cc
View File

@ -24,16 +24,7 @@ namespace {
using MslGeneratorImplTest = testing::Test;
TEST_F(MslGeneratorImplTest, DISABLED_EmitExpression_Identifier) {
ast::IdentifierExpression i(std::vector<std::string>{"std", "glsl"});
ast::Module m;
GeneratorImpl g(&m);
ASSERT_TRUE(g.EmitExpression(&i)) << g.error();
EXPECT_EQ(g.result(), "std::glsl");
}
TEST_F(MslGeneratorImplTest, EmitIdentifierExpression_Single) {
TEST_F(MslGeneratorImplTest, EmitIdentifierExpression) {
ast::IdentifierExpression i("foo");
ast::Module m;
@ -51,16 +42,6 @@ TEST_F(MslGeneratorImplTest, EmitIdentifierExpression_Single_WithCollision) {
EXPECT_EQ(g.result(), "virtual_tint_0");
}
// TODO(dsinclair): Handle import names
TEST_F(MslGeneratorImplTest, DISABLED_EmitIdentifierExpression_MultipleNames) {
ast::IdentifierExpression i({"std", "glsl", "init"});
ast::Module m;
GeneratorImpl g(&m);
ASSERT_TRUE(g.EmitExpression(&i)) << g.error();
EXPECT_EQ(g.result(), "std::glsl::init");
}
} // namespace
} // namespace msl
} // namespace writer

116
src/writer/msl/generator_impl_import_test.cc
View File

@ -57,46 +57,41 @@ TEST_P(MslImportData_SingleParamTest, FloatScalar) {
params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::FloatLiteral>(&f32, 1.f)));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", param.name}),
std::move(params));
auto ident = std::make_unique<ast::IdentifierExpression>(param.name);
auto* ident_ptr = ident.get();
ast::CallExpression call(std::move(ident), std::move(params));
Context ctx;
ast::Module mod;
TypeDeterminer td(&ctx, &mod);
mod.AddImport(std::make_unique<ast::Import>("GLSL.std.450", "std"));
ASSERT_TRUE(td.DetermineResultType(&expr)) << td.error();
// The call type determination will set the intrinsic data for the ident
ASSERT_TRUE(td.DetermineResultType(&call)) << td.error();
GeneratorImpl g(&mod);
ASSERT_TRUE(g.EmitImportFunction(&expr)) << g.error();
EXPECT_EQ(g.result(),
std::string("metal::") + param.msl_name + "(1.00000000f)");
ASSERT_TRUE(g.EmitBuiltinName(ident_ptr)) << g.error();
EXPECT_EQ(g.result(), std::string("metal::") + param.msl_name);
}
INSTANTIATE_TEST_SUITE_P(MslGeneratorImplTest,
MslImportData_SingleParamTest,
testing::Values(MslImportData{"acos", "acos"},
MslImportData{"acosh", "acosh"},
testing::Values(MslImportData{"abs", "fabs"},
MslImportData{"acos", "acos"},
MslImportData{"asin", "asin"},
MslImportData{"asinh", "asinh"},
MslImportData{"atan", "atan"},
MslImportData{"atanh", "atanh"},
MslImportData{"ceil", "ceil"},
MslImportData{"cos", "cos"},
MslImportData{"cosh", "cosh"},
MslImportData{"ceil", "ceil"},
MslImportData{"exp", "exp"},
MslImportData{"exp2", "exp2"},
MslImportData{"fabs", "fabs"},
MslImportData{"floor", "floor"},
MslImportData{"fract", "fract"},
MslImportData{"inversesqrt", "rsqrt"},
MslImportData{"inverseSqrt", "rsqrt"},
MslImportData{"length", "length"},
MslImportData{"log", "log"},
MslImportData{"log2", "log2"},
MslImportData{"normalize",
"normalize"},
MslImportData{"round", "round"},
MslImportData{"fsign", "sign"},
MslImportData{"sign", "sign"},
MslImportData{"sin", "sin"},
MslImportData{"sinh", "sinh"},
MslImportData{"sqrt", "sqrt"},
@ -111,19 +106,17 @@ TEST_F(MslGeneratorImplTest, MslImportData_SingleParamTest_IntScalar) {
params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::SintLiteral>(&i32, 1)));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", "sabs"}),
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>("abs"),
std::move(params));
Context ctx;
ast::Module mod;
TypeDeterminer td(&ctx, &mod);
mod.AddImport(std::make_unique<ast::Import>("GLSL.std.450", "std"));
ASSERT_TRUE(td.DetermineResultType(&expr)) << td.error();
GeneratorImpl g(&mod);
ASSERT_TRUE(g.EmitImportFunction(&expr)) << g.error();
ASSERT_TRUE(g.EmitCall(&expr)) << g.error();
EXPECT_EQ(g.result(), R"(metal::abs(1))");
}
@ -139,19 +132,18 @@ TEST_P(MslImportData_DualParamTest, FloatScalar) {
params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::FloatLiteral>(&f32, 2.f)));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", param.name}),
std::move(params));
ast::CallExpression expr(
std::make_unique<ast::IdentifierExpression>(param.name),
std::move(params));
Context ctx;
ast::Module mod;
TypeDeterminer td(&ctx, &mod);
mod.AddImport(std::make_unique<ast::Import>("GLSL.std.450", "std"));
ASSERT_TRUE(td.DetermineResultType(&expr)) << td.error();
GeneratorImpl g(&mod);
ASSERT_TRUE(g.EmitImportFunction(&expr)) << g.error();
ASSERT_TRUE(g.EmitCall(&expr)) << g.error();
EXPECT_EQ(g.result(), std::string("metal::") + param.msl_name +
"(1.00000000f, 2.00000000f)");
}
@ -159,10 +151,8 @@ INSTANTIATE_TEST_SUITE_P(MslGeneratorImplTest,
MslImportData_DualParamTest,
testing::Values(MslImportData{"atan2", "atan2"},
MslImportData{"distance", "distance"},
MslImportData{"fmax", "fmax"},
MslImportData{"fmin", "fmin"},
MslImportData{"nmax", "fmax"},
MslImportData{"nmin", "fmin"},
MslImportData{"max", "fmax"},
MslImportData{"min", "fmin"},
MslImportData{"pow", "pow"},
MslImportData{"reflect", "reflect"},
MslImportData{"step", "step"}));
@ -196,19 +186,18 @@ TEST_P(MslImportData_DualParam_VectorTest, FloatVector) {
params.push_back(std::make_unique<ast::TypeConstructorExpression>(
&vec, std::move(type_params)));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", param.name}),
std::move(params));
ast::CallExpression expr(
std::make_unique<ast::IdentifierExpression>(param.name),
std::move(params));
Context ctx;
ast::Module mod;
TypeDeterminer td(&ctx, &mod);
mod.AddImport(std::make_unique<ast::Import>("GLSL.std.450", "std"));
ASSERT_TRUE(td.DetermineResultType(&expr)) << td.error();
GeneratorImpl g(&mod);
ASSERT_TRUE(g.EmitImportFunction(&expr)) << g.error();
ASSERT_TRUE(g.EmitCall(&expr)) << g.error();
EXPECT_EQ(g.result(), std::string("metal::") + param.msl_name +
"(float3(1.00000000f, 2.00000000f, 3.00000000f), "
"float3(4.00000000f, 5.00000000f, 6.00000000f))");
@ -229,27 +218,24 @@ TEST_P(MslImportData_DualParam_Int_Test, IntScalar) {
params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::SintLiteral>(&i32, 2)));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", param.name}),
std::move(params));
ast::CallExpression expr(
std::make_unique<ast::IdentifierExpression>(param.name),
std::move(params));
Context ctx;
ast::Module mod;
TypeDeterminer td(&ctx, &mod);
mod.AddImport(std::make_unique<ast::Import>("GLSL.std.450", "std"));
ASSERT_TRUE(td.DetermineResultType(&expr)) << td.error();
GeneratorImpl g(&mod);
ASSERT_TRUE(g.EmitImportFunction(&expr)) << g.error();
ASSERT_TRUE(g.EmitCall(&expr)) << g.error();
EXPECT_EQ(g.result(), std::string("metal::") + param.msl_name + "(1, 2)");
}
INSTANTIATE_TEST_SUITE_P(MslGeneratorImplTest,
MslImportData_DualParam_Int_Test,
testing::Values(MslImportData{"smax", "max"},
MslImportData{"smin", "min"},
MslImportData{"umax", "max"},
MslImportData{"umin", "min"}));
testing::Values(MslImportData{"max", "max"},
MslImportData{"min", "min"}));
using MslImportData_TripleParamTest = testing::TestWithParam<MslImportData>;
TEST_P(MslImportData_TripleParamTest, FloatScalar) {
@ -265,31 +251,29 @@ TEST_P(MslImportData_TripleParamTest, FloatScalar) {
params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::FloatLiteral>(&f32, 3.f)));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", param.name}),
std::move(params));
ast::CallExpression expr(
std::make_unique<ast::IdentifierExpression>(param.name),
std::move(params));
Context ctx;
ast::Module mod;
TypeDeterminer td(&ctx, &mod);
mod.AddImport(std::make_unique<ast::Import>("GLSL.std.450", "std"));
ASSERT_TRUE(td.DetermineResultType(&expr)) << td.error();
GeneratorImpl g(&mod);
ASSERT_TRUE(g.EmitImportFunction(&expr)) << g.error();
ASSERT_TRUE(g.EmitCall(&expr)) << g.error();
EXPECT_EQ(g.result(), std::string("metal::") + param.msl_name +
"(1.00000000f, 2.00000000f, 3.00000000f)");
}
INSTANTIATE_TEST_SUITE_P(
MslGeneratorImplTest,
MslImportData_TripleParamTest,
testing::Values(MslImportData{"faceforward", "faceforward"},
testing::Values(MslImportData{"faceForward", "faceforward"},
MslImportData{"fma", "fma"},
MslImportData{"fclamp", "clamp"},
MslImportData{"fmix", "mix"},
MslImportData{"nclamp", "clamp"},
MslImportData{"smoothstep", "smoothstep"}));
MslImportData{"mix", "mix"},
MslImportData{"clamp", "clamp"},
MslImportData{"smoothStep", "smoothstep"}));
using MslImportData_TripleParam_Int_Test =
testing::TestWithParam<MslImportData>;
@ -306,25 +290,24 @@ TEST_P(MslImportData_TripleParam_Int_Test, IntScalar) {
params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::SintLiteral>(&i32, 3)));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", param.name}),
std::move(params));
ast::CallExpression expr(
std::make_unique<ast::IdentifierExpression>(param.name),
std::move(params));
Context ctx;
ast::Module mod;
TypeDeterminer td(&ctx, &mod);
mod.AddImport(std::make_unique<ast::Import>("GLSL.std.450", "std"));
ASSERT_TRUE(td.DetermineResultType(&expr)) << td.error();
GeneratorImpl g(&mod);
ASSERT_TRUE(g.EmitImportFunction(&expr)) << g.error();
ASSERT_TRUE(g.EmitCall(&expr)) << g.error();
EXPECT_EQ(g.result(), std::string("metal::") + param.msl_name + "(1, 2, 3)");
}
INSTANTIATE_TEST_SUITE_P(MslGeneratorImplTest,
MslImportData_TripleParam_Int_Test,
testing::Values(MslImportData{"sclamp", "clamp"},
MslImportData{"uclamp", "clamp"}));
testing::Values(MslImportData{"clamp", "clamp"},
MslImportData{"clamp", "clamp"}));
TEST_F(MslGeneratorImplTest, MslImportData_Determinant) {
ast::type::F32Type f32;
@ -336,14 +319,13 @@ TEST_F(MslGeneratorImplTest, MslImportData_Determinant) {
ast::ExpressionList params;
params.push_back(std::make_unique<ast::IdentifierExpression>("var"));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", "determinant"}),
std::move(params));
ast::CallExpression expr(
std::make_unique<ast::IdentifierExpression>("determinant"),
std::move(params));
Context ctx;
ast::Module mod;
mod.AddGlobalVariable(std::move(var));
mod.AddImport(std::make_unique<ast::Import>("GLSL.std.450", "std"));
TypeDeterminer td(&ctx, &mod);
// Register the global
@ -351,7 +333,7 @@ TEST_F(MslGeneratorImplTest, MslImportData_Determinant) {
ASSERT_TRUE(td.DetermineResultType(&expr)) << td.error();
GeneratorImpl g(&mod);
ASSERT_TRUE(g.EmitImportFunction(&expr)) << g.error();
ASSERT_TRUE(g.EmitCall(&expr)) << g.error();
EXPECT_EQ(g.result(), std::string("metal::determinant(var)"));
}

52
src/writer/msl/generator_impl_intrinsic_test.cc
View File

@ -30,11 +30,11 @@ namespace {
using MslGeneratorImplTest = testing::Test;
struct IntrinsicData {
const char* name;
ast::Intrinsic intrinsic;
const char* msl_name;
};
inline std::ostream& operator<<(std::ostream& out, IntrinsicData data) {
out << data.name;
out << data.msl_name;
return out;
}
using MslIntrinsicTest = testing::TestWithParam<IntrinsicData>;
@ -43,28 +43,28 @@ TEST_P(MslIntrinsicTest, Emit) {
ast::Module m;
GeneratorImpl g(&m);
EXPECT_EQ(g.generate_intrinsic_name(param.name), param.msl_name);
EXPECT_EQ(g.generate_intrinsic_name(param.intrinsic), param.msl_name);
}
INSTANTIATE_TEST_SUITE_P(MslGeneratorImplTest,
MslIntrinsicTest,
testing::Values(IntrinsicData{"any", "any"},
IntrinsicData{"all", "all"},
IntrinsicData{"dot", "dot"},
IntrinsicData{"dpdx", "dfdx"},
IntrinsicData{"dpdx_coarse", "dfdx"},
IntrinsicData{"dpdx_fine", "dfdx"},
IntrinsicData{"dpdy", "dfdy"},
IntrinsicData{"dpdy_coarse", "dfdy"},
IntrinsicData{"dpdy_fine", "dfdy"},
IntrinsicData{"fwidth", "fwidth"},
IntrinsicData{"fwidth_coarse",
"fwidth"},
IntrinsicData{"fwidth_fine", "fwidth"},
IntrinsicData{"is_finite", "isfinite"},
IntrinsicData{"is_inf", "isinf"},
IntrinsicData{"is_nan", "isnan"},
IntrinsicData{"is_normal", "isnormal"},
IntrinsicData{"select", "select"}));
INSTANTIATE_TEST_SUITE_P(
MslGeneratorImplTest,
MslIntrinsicTest,
testing::Values(IntrinsicData{ast::Intrinsic::kAny, "any"},
IntrinsicData{ast::Intrinsic::kAll, "all"},
IntrinsicData{ast::Intrinsic::kDot, "dot"},
IntrinsicData{ast::Intrinsic::kDpdx, "dfdx"},
IntrinsicData{ast::Intrinsic::kDpdxCoarse, "dfdx"},
IntrinsicData{ast::Intrinsic::kDpdxFine, "dfdx"},
IntrinsicData{ast::Intrinsic::kDpdy, "dfdy"},
IntrinsicData{ast::Intrinsic::kDpdyCoarse, "dfdy"},
IntrinsicData{ast::Intrinsic::kDpdyFine, "dfdy"},
IntrinsicData{ast::Intrinsic::kFwidth, "fwidth"},
IntrinsicData{ast::Intrinsic::kFwidthCoarse, "fwidth"},
IntrinsicData{ast::Intrinsic::kFwidthFine, "fwidth"},
IntrinsicData{ast::Intrinsic::kIsFinite, "isfinite"},
IntrinsicData{ast::Intrinsic::kIsInf, "isinf"},
IntrinsicData{ast::Intrinsic::kIsNan, "isnan"},
IntrinsicData{ast::Intrinsic::kIsNormal, "isnormal"},
IntrinsicData{ast::Intrinsic::kSelect, "select"}));
TEST_F(MslGeneratorImplTest, DISABLED_Intrinsic_OuterProduct) {
ast::type::F32Type f32;
@ -106,7 +106,7 @@ TEST_F(MslGeneratorImplTest, DISABLED_Intrinsic_OuterProduct) {
TEST_F(MslGeneratorImplTest, Intrinsic_Bad_Name) {
ast::Module m;
GeneratorImpl g(&m);
EXPECT_EQ(g.generate_intrinsic_name("unknown name"), "");
EXPECT_EQ(g.generate_intrinsic_name(ast::Intrinsic::kNone), "");
}
TEST_F(MslGeneratorImplTest, Intrinsic_Call) {
@ -117,7 +117,11 @@ TEST_F(MslGeneratorImplTest, Intrinsic_Call) {
ast::CallExpression call(std::make_unique<ast::IdentifierExpression>("dot"),
std::move(params));
Context ctx;
ast::Module m;
TypeDeterminer td(&ctx, &m);
ASSERT_TRUE(td.DetermineResultType(&call)) << td.error();
GeneratorImpl g(&m);
g.increment_indent();
ASSERT_TRUE(g.EmitExpression(&call)) << g.error();

324
src/writer/spirv/builder.cc
View File

@ -17,6 +17,7 @@
#include <sstream>
#include <utility>
#include "spirv/unified1/GLSL.std.450.h"
#include "spirv/unified1/spirv.h"
#include "src/ast/array_accessor_expression.h"
#include "src/ast/as_expression.h"
@ -68,11 +69,15 @@
#include "src/ast/unary_op_expression.h"
#include "src/ast/variable_decl_statement.h"
#include <iostream>
namespace tint {
namespace writer {
namespace spirv {
namespace {
const char kGLSLstd450[] = "GLSL.std.450";
uint32_t size_of(const InstructionList& instructions) {
uint32_t size = 0;
for (const auto& inst : instructions)
@ -147,6 +152,119 @@ ast::type::MatrixType* GetNestedMatrixType(ast::type::Type* type) {
return type->IsMatrix() ? type->AsMatrix() : nullptr;
}
uint32_t intrinsic_to_glsl_method(ast::type::Type* type,
ast::Intrinsic intrinsic) {
switch (intrinsic) {
case ast::Intrinsic::kAbs:
if (type->is_float_scalar_or_vector()) {
return GLSLstd450FAbs;
} else {
return GLSLstd450SAbs;
}
case ast::Intrinsic::kAcos:
return GLSLstd450Acos;
case ast::Intrinsic::kAsin:
return GLSLstd450Asin;
case ast::Intrinsic::kAtan:
return GLSLstd450Atan;
case ast::Intrinsic::kAtan2:
return GLSLstd450Atan2;
case ast::Intrinsic::kCeil:
return GLSLstd450Ceil;
case ast::Intrinsic::kClamp:
if (type->is_float_scalar_or_vector()) {
return GLSLstd450NClamp;
} else if (type->is_unsigned_scalar_or_vector()) {
return GLSLstd450UClamp;
} else {
return GLSLstd450SClamp;
}
case ast::Intrinsic::kCos:
return GLSLstd450Cos;
case ast::Intrinsic::kCosh:
return GLSLstd450Cosh;
case ast::Intrinsic::kCross:
return GLSLstd450Cross;
case ast::Intrinsic::kDeterminant:
return GLSLstd450Determinant;
case ast::Intrinsic::kDistance:
return GLSLstd450Distance;
case ast::Intrinsic::kExp:
return GLSLstd450Exp;
case ast::Intrinsic::kExp2:
return GLSLstd450Exp2;
case ast::Intrinsic::kFaceForward:
return GLSLstd450FaceForward;
case ast::Intrinsic::kFloor:
return GLSLstd450Floor;
case ast::Intrinsic::kFma:
return GLSLstd450Fma;
case ast::Intrinsic::kFract:
return GLSLstd450Fract;
case ast::Intrinsic::kFrexp:
return GLSLstd450Frexp;
case ast::Intrinsic::kInverseSqrt:
return GLSLstd450InverseSqrt;
case ast::Intrinsic::kLdexp:
return GLSLstd450Ldexp;
case ast::Intrinsic::kLength:
return GLSLstd450Length;
case ast::Intrinsic::kLog:
return GLSLstd450Log;
case ast::Intrinsic::kLog2:
return GLSLstd450Log2;
case ast::Intrinsic::kMax:
if (type->is_float_scalar_or_vector()) {
return GLSLstd450NMax;
} else if (type->is_unsigned_scalar_or_vector()) {
return GLSLstd450UMax;
} else {
return GLSLstd450SMax;
}
case ast::Intrinsic::kMin:
if (type->is_float_scalar_or_vector()) {
return GLSLstd450NMin;
} else if (type->is_unsigned_scalar_or_vector()) {
return GLSLstd450UMin;
} else {
return GLSLstd450SMin;
}
case ast::Intrinsic::kMix:
return GLSLstd450FMix;
case ast::Intrinsic::kModf:
return GLSLstd450Modf;
case ast::Intrinsic::kNormalize:
return GLSLstd450Normalize;
case ast::Intrinsic::kPow:
return GLSLstd450Pow;
case ast::Intrinsic::kReflect:
return GLSLstd450Reflect;
case ast::Intrinsic::kRound:
return GLSLstd450Round;
case ast::Intrinsic::kSign:
return GLSLstd450FSign;
case ast::Intrinsic::kSin:
return GLSLstd450Sin;
case ast::Intrinsic::kSinh:
return GLSLstd450Sinh;
case ast::Intrinsic::kSmoothStep:
return GLSLstd450SmoothStep;
case ast::Intrinsic::kSqrt:
return GLSLstd450Sqrt;
case ast::Intrinsic::kStep:
return GLSLstd450Step;
case ast::Intrinsic::kTan:
return GLSLstd450Tan;
case ast::Intrinsic::kTanh:
return GLSLstd450Tanh;
case ast::Intrinsic::kTrunc:
return GLSLstd450Trunc;
default:
break;
}
return 0;
}
} // namespace
Builder::AccessorInfo::AccessorInfo() : source_id(0), source_type(nullptr) {}
@ -165,10 +283,6 @@ bool Builder::Build() {
push_preamble(spv::Op::OpExtension,
{Operand::String("SPV_KHR_vulkan_memory_model")});
for (const auto& imp : mod_->imports()) {
GenerateImport(imp.get());
}
push_preamble(spv::Op::OpMemoryModel,
{Operand::Int(SpvAddressingModelLogical),
Operand::Int(SpvMemoryModelVulkanKHR)});
@ -853,18 +967,6 @@ uint32_t Builder::GenerateAccessorExpression(ast::Expression* expr) {
uint32_t Builder::GenerateIdentifierExpression(
ast::IdentifierExpression* expr) {
uint32_t val = 0;
if (expr->has_path()) {
auto* imp = mod_->FindImportByName(expr->path());
if (imp == nullptr) {
error_ = "unable to find import for " + expr->path();
return 0;
}
val = imp->GetIdForMethod(expr->name());
if (val == 0) {
error_ = "unable to lookup: " + expr->name() + " in " + expr->path();
}
return val;
}
if (scope_stack_.get(expr->name(), &val)) {
return val;
}
@ -920,14 +1022,18 @@ uint32_t Builder::GenerateUnaryOpExpression(ast::UnaryOpExpression* expr) {
return result_id;
}
void Builder::GenerateImport(ast::Import* imp) {
void Builder::GenerateGLSLstd450Import() {
if (import_name_to_id_.find(kGLSLstd450) != import_name_to_id_.end()) {
return;
}
auto result = result_op();
auto id = result.to_i();
push_preamble(spv::Op::OpExtInstImport,
{result, Operand::String(imp->path())});
{result, Operand::String(kGLSLstd450)});
import_name_to_id_[imp->name()] = id;
import_name_to_id_[kGLSLstd450] = id;
}
uint32_t Builder::GenerateConstructorExpression(
@ -1361,8 +1467,8 @@ uint32_t Builder::GenerateCallExpression(ast::CallExpression* expr) {
auto* ident = expr->func()->AsIdentifier();
if (!ident->has_path() && ast::intrinsic::IsIntrinsic(ident->name())) {
return GenerateIntrinsic(ident->name(), expr);
if (ident->IsIntrinsic()) {
return GenerateIntrinsic(ident, expr);
}
auto type_id = GenerateTypeIfNeeded(expr->func()->result_type());
@ -1373,45 +1479,14 @@ uint32_t Builder::GenerateCallExpression(ast::CallExpression* expr) {
auto result = result_op();
auto result_id = result.to_i();
spv::Op op = spv::Op::OpNop;
OperandList ops = {Operand::Int(type_id), result};
// Handle regular function calls
if (!ident->has_path()) {
auto func_id = func_name_to_id_[ident->name()];
if (func_id == 0) {
error_ = "unable to find called function: " + ident->name();
return 0;
}
ops.push_back(Operand::Int(func_id));
op = spv::Op::OpFunctionCall;
} else {
// Imported function call
auto set_iter = import_name_to_id_.find(ident->path());
if (set_iter == import_name_to_id_.end()) {
error_ = "unknown import " + ident->path();
return 0;
}
auto set_id = set_iter->second;
auto* imp = mod_->FindImportByName(ident->path());
if (imp == nullptr) {
error_ = "unknown import " + ident->path();
return 0;
}
auto inst_id = imp->GetIdForMethod(ident->name());
if (inst_id == 0) {
error_ = "unknown method " + ident->name();
return 0;
}
ops.push_back(Operand::Int(set_id));
ops.push_back(Operand::Int(inst_id));
op = spv::Op::OpExtInst;
auto func_id = func_name_to_id_[ident->name()];
if (func_id == 0) {
error_ = "unable to find called function: " + ident->name();
return 0;
}
ops.push_back(Operand::Int(func_id));
for (const auto& param : expr->params()) {
auto id = GenerateExpression(param.get());
@ -1422,12 +1497,12 @@ uint32_t Builder::GenerateCallExpression(ast::CallExpression* expr) {
ops.push_back(Operand::Int(id));
}
push_function_inst(op, std::move(ops));
push_function_inst(spv::Op::OpFunctionCall, std::move(ops));
return result_id;
}
uint32_t Builder::GenerateIntrinsic(const std::string& name,
uint32_t Builder::GenerateIntrinsic(ast::IdentifierExpression* ident,
ast::CallExpression* call) {
auto result = result_op();
auto result_id = result.to_i();
@ -1438,6 +1513,76 @@ uint32_t Builder::GenerateIntrinsic(const std::string& name,
}
OperandList params = {Operand::Int(result_type_id), result};
auto intrinsic = ident->intrinsic();
if (ast::intrinsic::IsFineDerivative(intrinsic) ||
ast::intrinsic::IsCoarseDerivative(intrinsic)) {
push_capability(SpvCapabilityDerivativeControl);
}
spv::Op op = spv::Op::OpNop;
if (intrinsic == ast::Intrinsic::kAny) {
op = spv::Op::OpAny;
} else if (intrinsic == ast::Intrinsic::kAll) {
op = spv::Op::OpAll;
} else if (intrinsic == ast::Intrinsic::kCountOneBits) {
op = spv::Op::OpBitCount;
} else if (intrinsic == ast::Intrinsic::kDot) {
op = spv::Op::OpDot;
} else if (intrinsic == ast::Intrinsic::kDpdx) {
op = spv::Op::OpDPdx;
} else if (intrinsic == ast::Intrinsic::kDpdxCoarse) {
op = spv::Op::OpDPdxCoarse;
} else if (intrinsic == ast::Intrinsic::kDpdxFine) {
op = spv::Op::OpDPdxFine;
} else if (intrinsic == ast::Intrinsic::kDpdy) {
op = spv::Op::OpDPdy;
} else if (intrinsic == ast::Intrinsic::kDpdyCoarse) {
op = spv::Op::OpDPdyCoarse;
} else if (intrinsic == ast::Intrinsic::kDpdyFine) {
op = spv::Op::OpDPdyFine;
} else if (intrinsic == ast::Intrinsic::kFwidth) {
op = spv::Op::OpFwidth;
} else if (intrinsic == ast::Intrinsic::kFwidthCoarse) {
op = spv::Op::OpFwidthCoarse;
} else if (intrinsic == ast::Intrinsic::kFwidthFine) {
op = spv::Op::OpFwidthFine;
} else if (intrinsic == ast::Intrinsic::kIsInf) {
op = spv::Op::OpIsInf;
} else if (intrinsic == ast::Intrinsic::kIsNan) {
op = spv::Op::OpIsNan;
} else if (intrinsic == ast::Intrinsic::kOuterProduct) {
op = spv::Op::OpOuterProduct;
} else if (intrinsic == ast::Intrinsic::kReverseBits) {
op = spv::Op::OpBitReverse;
} else if (intrinsic == ast::Intrinsic::kSelect) {
op = spv::Op::OpSelect;
} else if (!ast::intrinsic::IsTextureIntrinsic(intrinsic)) {
GenerateGLSLstd450Import();
auto set_iter = import_name_to_id_.find(kGLSLstd450);
if (set_iter == import_name_to_id_.end()) {
error_ = std::string("unknown import ") + kGLSLstd450;
return 0;
}
auto set_id = set_iter->second;
auto inst_id =
intrinsic_to_glsl_method(ident->result_type(), ident->intrinsic());
if (inst_id == 0) {
error_ = "unknown method " + ident->name();
return 0;
}
params.push_back(Operand::Int(set_id));
params.push_back(Operand::Int(inst_id));
op = spv::Op::OpExtInst;
}
if (!ast::intrinsic::IsTextureIntrinsic(intrinsic) && op == spv::Op::OpNop) {
error_ = "unable to determine operator for: " + ident->name();
return 0;
}
for (const auto& p : call->params()) {
auto val_id = GenerateExpression(p.get());
if (val_id == 0) {
@ -1448,59 +1593,16 @@ uint32_t Builder::GenerateIntrinsic(const std::string& name,
params.push_back(Operand::Int(val_id));
}
if (ast::intrinsic::IsTextureOperationIntrinsic(name)) {
return GenerateTextureIntrinsic(name, call, result_id, params);
if (ast::intrinsic::IsTextureIntrinsic(intrinsic)) {
return GenerateTextureIntrinsic(ident, call, result_id, params);
}
if (ast::intrinsic::IsFineDerivative(name) ||
ast::intrinsic::IsCoarseDerivative(name)) {
push_capability(SpvCapabilityDerivativeControl);
}
spv::Op op = spv::Op::OpNop;
if (name == "any") {
op = spv::Op::OpAny;
} else if (name == "all") {
op = spv::Op::OpAll;
} else if (name == "dot") {
op = spv::Op::OpDot;
} else if (name == "dpdx") {
op = spv::Op::OpDPdx;
} else if (name == "dpdxCoarse") {
op = spv::Op::OpDPdxCoarse;
} else if (name == "dpdxFine") {
op = spv::Op::OpDPdxFine;
} else if (name == "dpdy") {
op = spv::Op::OpDPdy;
} else if (name == "dpdyCoarse") {
op = spv::Op::OpDPdyCoarse;
} else if (name == "dpdyFine") {
op = spv::Op::OpDPdyFine;
} else if (name == "fwidth") {
op = spv::Op::OpFwidth;
} else if (name == "fwidthCoarse") {
op = spv::Op::OpFwidthCoarse;
} else if (name == "fwidthFine") {
op = spv::Op::OpFwidthFine;
} else if (name == "isInf") {
op = spv::Op::OpIsInf;
} else if (name == "isNan") {
op = spv::Op::OpIsNan;
} else if (name == "outerProduct") {
op = spv::Op::OpOuterProduct;
} else if (name == "select") {
op = spv::Op::OpSelect;
}
if (op == spv::Op::OpNop) {
error_ = "unable to determine operator for: " + name;
return 0;
}
push_function_inst(op, params);
return result_id;
}
uint32_t Builder::GenerateTextureIntrinsic(const std::string& name,
uint32_t Builder::GenerateTextureIntrinsic(ast::IdentifierExpression* ident,
ast::CallExpression* call,
uint32_t result_id,
OperandList wgsl_params) {
@ -1512,7 +1614,7 @@ uint32_t Builder::GenerateTextureIntrinsic(const std::string& name,
// TODO: Remove the LOD param from textureLoad on storage textures when
// https://github.com/gpuweb/gpuweb/pull/1032 gets merged.
if (name == "textureLoad") {
if (ident->intrinsic() == ast::Intrinsic::kTextureLoad) {
std::vector<Operand> spirv_params = {
std::move(wgsl_params[0]), std::move(wgsl_params[1]),
std::move(wgsl_params[2]), std::move(wgsl_params[3])};
@ -1538,17 +1640,17 @@ uint32_t Builder::GenerateTextureIntrinsic(const std::string& name,
std::move(wgsl_params[3]))),
std::move(wgsl_params[4])};
if (name == "textureSample") {
if (ident->intrinsic() == ast::Intrinsic::kTextureSample) {
op = spv::Op::OpImageSampleImplicitLod;
} else if (name == "textureSampleLevel") {
} else if (ident->intrinsic() == ast::Intrinsic::kTextureSampleLevel) {
op = spv::Op::OpImageSampleExplicitLod;
spirv_params.push_back(Operand::Int(SpvImageOperandsLodMask));
spirv_params.push_back(std::move(wgsl_params[5]));
} else if (name == "textureSampleBias") {
} else if (ident->intrinsic() == ast::Intrinsic::kTextureSampleBias) {
op = spv::Op::OpImageSampleImplicitLod;
spirv_params.push_back(Operand::Int(SpvImageOperandsBiasMask));
spirv_params.push_back(std::move(wgsl_params[5]));
} else if (name == "textureSampleCompare") {
} else if (ident->intrinsic() == ast::Intrinsic::kTextureSampleCompare) {
op = spv::Op::OpImageSampleDrefExplicitLod;
spirv_params.push_back(std::move(wgsl_params[5]));
@ -1558,7 +1660,7 @@ uint32_t Builder::GenerateTextureIntrinsic(const std::string& name,
spirv_params.push_back(Operand::Int(GenerateLiteralIfNeeded(&float_0)));
}
if (op == spv::Op::OpNop) {
error_ = "unable to determine operator for: " + name;
error_ = "unable to determine operator for: " + ident->name();
return 0;
}
push_function_inst(op, spirv_params);

11
src/writer/spirv/builder.h
View File

@ -247,8 +247,7 @@ class Builder {
/// @returns true on success
bool GenerateIfStatement(ast::IfStatement* stmt);
/// Generates an import instruction
/// @param imp the import
void GenerateImport(ast::Import* imp);
void GenerateGLSLstd450Import();
/// Generates a constructor expression
/// @param expr the expression to generate
/// @param is_global_init set true if this is a global variable constructor
@ -283,19 +282,19 @@ class Builder {
/// @returns the expression ID on success or 0 otherwise
uint32_t GenerateCallExpression(ast::CallExpression* expr);
/// Generates an intrinsic call
/// @param name the intrinsic name
/// @param ident the intrinsic expression
/// @param call the call expression
/// @returns the expression ID on success or 0 otherwise
uint32_t GenerateIntrinsic(const std::string& name,
uint32_t GenerateIntrinsic(ast::IdentifierExpression* ident,
ast::CallExpression* call);
/// Generates a texture intrinsic call
/// @param name the texture intrinsic name
/// @param ident the texture intrinsic
/// @param call the call expression
/// @param result_id result ID of the texture instruction
/// @param wgsl_params SPIR-V arguments for WGSL-specific intrinsic's call
/// parameters
/// @returns the expression ID on success or 0 otherwise
uint32_t GenerateTextureIntrinsic(const std::string& name,
uint32_t GenerateTextureIntrinsic(ast::IdentifierExpression* ident,
ast::CallExpression* call,
uint32_t result_id,
OperandList wgsl_params);

92
src/writer/spirv/builder_call_test.cc
View File

@ -38,98 +38,6 @@ namespace {
using BuilderTest = testing::Test;
TEST_F(BuilderTest, Call_GLSLMethod) {
ast::type::F32Type f32;
ast::type::VoidType void_type;
ast::ExpressionList params;
params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::FloatLiteral>(&f32, 1.f)));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", "round"}),
std::move(params));
Context ctx;
ast::Module mod;
TypeDeterminer td(&ctx, &mod);
auto imp = std::make_unique<ast::Import>("GLSL.std.450", "std");
auto* glsl = imp.get();
mod.AddImport(std::move(imp));
ASSERT_TRUE(td.DetermineResultType(&expr)) << td.error();
ast::Function func("a_func", {}, &void_type);
Builder b(&mod);
b.GenerateImport(glsl);
ASSERT_TRUE(b.GenerateFunction(&func)) << b.error();
EXPECT_EQ(b.GenerateCallExpression(&expr), 7u) << b.error();
EXPECT_EQ(DumpBuilder(b), R"(%1 = OpExtInstImport "GLSL.std.450"
OpName %4 "a_func"
%3 = OpTypeVoid
%2 = OpTypeFunction %3
%6 = OpTypeFloat 32
%8 = OpConstant %6 1
%4 = OpFunction %3 None %2
%5 = OpLabel
%7 = OpExtInst %6 %1 Round %8
OpFunctionEnd
)");
}
TEST_F(BuilderTest, Call_GLSLMethod_WithLoad) {
ast::type::F32Type f32;
ast::type::VoidType void_type;
auto var = std::make_unique<ast::Variable>("ident",
ast::StorageClass::kPrivate, &f32);
ast::ExpressionList params;
params.push_back(std::make_unique<ast::IdentifierExpression>("ident"));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", "round"}),
std::move(params));
Context ctx;
ast::Module mod;
TypeDeterminer td(&ctx, &mod);
td.RegisterVariableForTesting(var.get());
auto imp = std::make_unique<ast::Import>("GLSL.std.450", "std");
auto* glsl = imp.get();
mod.AddImport(std::move(imp));
ASSERT_TRUE(td.DetermineResultType(&expr)) << td.error();
ast::Function func("a_func", {}, &void_type);
Builder b(&mod);
b.GenerateImport(glsl);
ASSERT_TRUE(b.GenerateGlobalVariable(var.get())) << b.error();
ASSERT_TRUE(b.GenerateFunction(&func)) << b.error();
EXPECT_EQ(b.GenerateCallExpression(&expr), 10u) << b.error();
EXPECT_EQ(DumpBuilder(b), R"(%1 = OpExtInstImport "GLSL.std.450"
OpName %2 "ident"
OpName %8 "a_func"
%4 = OpTypeFloat 32
%3 = OpTypePointer Private %4
%5 = OpConstantNull %4
%2 = OpVariable %3 Private %5
%7 = OpTypeVoid
%6 = OpTypeFunction %7
%8 = OpFunction %7 None %6
%9 = OpLabel
%11 = OpLoad %4 %2
%10 = OpExtInst %4 %1 Round %11
OpFunctionEnd
)");
}
TEST_F(BuilderTest, Expression_Call) {
ast::type::F32Type f32;
ast::type::VoidType void_type;

1501
src/writer/spirv/builder_intrinsic_test.cc
View File

File diff suppressed because it is too large Load Diff

20
src/writer/spirv/builder_test.cc
View File

@ -19,7 +19,6 @@
#include "gtest/gtest.h"
#include "spirv/unified1/spirv.h"
#include "spirv/unified1/spirv.hpp11"
#include "src/ast/import.h"
#include "src/ast/module.h"
#include "src/writer/spirv/spv_dump.h"
@ -30,24 +29,7 @@ namespace {
using BuilderTest = testing::Test;
TEST_F(BuilderTest, InsertsPreambleWithImport) {
ast::Module m;
m.AddImport(std::make_unique<ast::Import>("GLSL.std.450", "glsl"));
Builder b(&m);
ASSERT_TRUE(b.Build());
ASSERT_EQ(b.capabilities().size(), 2u);
ASSERT_EQ(b.preamble().size(), 3u);
EXPECT_EQ(DumpBuilder(b), R"(OpCapability Shader
OpCapability VulkanMemoryModel
OpExtension "SPV_KHR_vulkan_memory_model"
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical Vulkan
)");
}
TEST_F(BuilderTest, InsertsPreambleWithoutImport) {
TEST_F(BuilderTest, InsertsPreamble) {
ast::Module m;
Builder b(&m);
ASSERT_TRUE(b.Build());

29
src/writer/wgsl/generator_impl.cc
View File

@ -75,15 +75,6 @@ GeneratorImpl::GeneratorImpl() = default;
GeneratorImpl::~GeneratorImpl() = default;
bool GeneratorImpl::Generate(const ast::Module& module) {
for (const auto& import : module.imports()) {
if (!EmitImport(import.get())) {
return false;
}
}
if (!module.imports().empty()) {
out_ << std::endl;
}
for (auto* const alias : module.alias_types()) {
if (!EmitAliasType(alias)) {
return false;
@ -114,16 +105,6 @@ bool GeneratorImpl::Generate(const ast::Module& module) {
bool GeneratorImpl::GenerateEntryPoint(const ast::Module& module,
ast::PipelineStage stage,
const std::string& name) {
// TODO(dsinclair): We're always emitting imports even if they aren't needed.
for (const auto& import : module.imports()) {
if (!EmitImport(import.get())) {
return false;
}
}
if (!module.imports().empty()) {
out_ << std::endl;
}
auto* func = module.FindFunctionByNameAndStage(name, stage);
if (func == nullptr) {
error_ = "Unable to find requested entry point: " + name;
@ -365,20 +346,10 @@ bool GeneratorImpl::EmitLiteral(ast::Literal* lit) {
bool GeneratorImpl::EmitIdentifier(ast::IdentifierExpression* expr) {
auto* ident = expr->AsIdentifier();
if (ident->has_path()) {
out_ << ident->path() << "::";
}
out_ << ident->name();
return true;
}
bool GeneratorImpl::EmitImport(const ast::Import* import) {
make_indent();
out_ << R"(import ")" << import->path() << R"(" as )" << import->name() << ";"
<< std::endl;
return true;
}
bool GeneratorImpl::EmitFunction(ast::Function* func) {
for (auto& deco : func->decorations()) {
make_indent();

5
src/writer/wgsl/generator_impl.h
View File

@ -21,7 +21,6 @@
#include "src/ast/array_accessor_expression.h"
#include "src/ast/constructor_expression.h"
#include "src/ast/identifier_expression.h"
#include "src/ast/import.h"
#include "src/ast/module.h"
#include "src/ast/scalar_constructor_expression.h"
#include "src/ast/type/alias_type.h"
@ -135,10 +134,6 @@ class GeneratorImpl : public TextGenerator {
/// @param stmt the statement to emit
/// @returns true if the statement was successfully emitted
bool EmitIf(ast::IfStatement* stmt);
/// Handles generating an import command
/// @param import the import to generate
/// @returns true if the import was emitted
bool EmitImport(const ast::Import* import);
/// Handles generating constructor expressions
/// @param expr the constructor expression
/// @returns true if the expression was emitted

16
src/writer/wgsl/generator_impl_identifier_test.cc
View File

@ -22,14 +22,6 @@ namespace {
using WgslGeneratorImplTest = testing::Test;
TEST_F(WgslGeneratorImplTest, EmitExpression_Identifier) {
ast::IdentifierExpression i(std::vector<std::string>{"std", "glsl"});
GeneratorImpl g;
ASSERT_TRUE(g.EmitExpression(&i)) << g.error();
EXPECT_EQ(g.result(), "std::glsl");
}
TEST_F(WgslGeneratorImplTest, EmitIdentifierExpression_Single) {
ast::IdentifierExpression i("glsl");
@ -38,14 +30,6 @@ TEST_F(WgslGeneratorImplTest, EmitIdentifierExpression_Single) {
EXPECT_EQ(g.result(), "glsl");
}
TEST_F(WgslGeneratorImplTest, EmitIdentifierExpression_MultipleNames) {
ast::IdentifierExpression i({"std", "glsl", "init"});
GeneratorImpl g;
ASSERT_TRUE(g.EmitExpression(&i)) << g.error();
EXPECT_EQ(g.result(), "std::glsl::init");
}
} // namespace
} // namespace wgsl
} // namespace writer

37
src/writer/wgsl/generator_impl_import_test.cc
View File

@ -1,37 +0,0 @@
// Copyright 2020 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include "gtest/gtest.h"
#include "src/writer/wgsl/generator_impl.h"
namespace tint {
namespace writer {
namespace wgsl {
namespace {
using WgslGeneratorImplTest = testing::Test;
TEST_F(WgslGeneratorImplTest, EmitImport) {
ast::Import import("GLSL.std.450", "std::glsl");
GeneratorImpl g;
ASSERT_TRUE(g.EmitImport(&import)) << g.error();
EXPECT_EQ(g.result(), R"(import "GLSL.std.450" as std::glsl;
)");
}
} // namespace
} // namespace wgsl
} // namespace writer
} // namespace tint

11
src/writer/wgsl/generator_impl_test.cc
View File

@ -17,6 +17,9 @@
#include <memory>
#include "gtest/gtest.h"
#include "src/ast/function.h"
#include "src/ast/module.h"
#include "src/ast/type/void_type.h"
namespace tint {
namespace writer {
@ -26,13 +29,17 @@ namespace {
using WgslGeneratorImplTest = testing::Test;
TEST_F(WgslGeneratorImplTest, Generate) {
ast::type::VoidType void_type;
ast::Module m;
m.AddImport(std::make_unique<ast::Import>("GLSL.std.430", "a"));
m.AddFunction(std::make_unique<ast::Function>("my_func", ast::VariableList{},
&void_type));
GeneratorImpl g;
ASSERT_TRUE(g.Generate(m)) << g.error();
EXPECT_EQ(g.result(), R"(import "GLSL.std.430" as a;
EXPECT_EQ(g.result(), R"(fn my_func() -> void {
}
)");
}

16
test/compute_boids.wgsl
View File

@ -12,8 +12,6 @@
# See the License for the specific language governing permissions and
# limitations under the License.
import "GLSL.std.450" as std;
# vertex shader
[[location(0)]] var<in> a_particlePos : vec2<f32>;
@ -23,10 +21,10 @@ import "GLSL.std.450" as std;
[[stage(vertex)]]
fn vert_main() -> void {
var angle : f32 = -std::atan2(a_particleVel.x, a_particleVel.y);
var angle : f32 = -atan2(a_particleVel.x, a_particleVel.y);
var pos : vec2<f32> = vec2<f32>(
(a_pos.x * std::cos(angle)) - (a_pos.y * std::sin(angle)),
(a_pos.x * std::sin(angle)) + (a_pos.y * std::cos(angle)));
(a_pos.x * cos(angle)) - (a_pos.y * sin(angle)),
(a_pos.x * sin(angle)) + (a_pos.y * cos(angle)));
gl_Position = vec4<f32>(pos + a_particlePos, 0.0, 1.0);
return;
}
@ -97,14 +95,14 @@ fn comp_main() -> void {
pos = particlesA.particles[i].pos.xy;
vel = particlesA.particles[i].vel.xy;
if (std::distance(pos, vPos) < params.rule1Distance) {
if (distance(pos, vPos) < params.rule1Distance) {
cMass = cMass + pos;
cMassCount = cMassCount + 1;
}
if (std::distance(pos, vPos) < params.rule2Distance) {
if (distance(pos, vPos) < params.rule2Distance) {
colVel = colVel - (pos - vPos);
}
if (std::distance(pos, vPos) < params.rule3Distance) {
if (distance(pos, vPos) < params.rule3Distance) {
cVel = cVel + vel;
cVelCount = cVelCount + 1;
}
@ -125,7 +123,7 @@ fn comp_main() -> void {
(cVel * params.rule3Scale);
# clamp velocity for a more pleasing simulation
vVel = std::normalize(vVel) * std::fclamp(std::length(vVel), 0.0, 0.1);
vVel = normalize(vVel) * clamp(length(vVel), 0.0, 0.1);
# kinematic update
vPos = vPos + (vVel * params.deltaT);