encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

Implement data packing intrinsics 

* Fix how the HLSL writer determines how to use a RWByteAddressBuffer
* Fix how the HLSL writer decides the register space for a storage
  variable
* Fix inference of hlsl format in the tint executable
* Add support for data packing intrinsics
  * type determination
  * validation
  * writers
  * spirv reader

Bug: tint:340, tint:473, tint:474
Change-Id: I45dc8fd7c6f9abc7d30f617c7e3d713d7965b76e
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/40342
Commit-Queue: Alan Baker <alanbaker@google.com>
Reviewed-by: dan sinclair <dsinclair@chromium.org>
This commit is contained in:
Alan Baker 2021-02-04 16:17:49 +00:00 committed by Commit Bot service account
parent fbd47c752e
commit c63e1c0791
16 changed files with 706 additions and 130 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

8
samples/main.cc
View File

@ -164,7 +164,13 @@ Format infer_format(const std::string& filename) {
if (ends_with(filename, ".metal")) {
return Format::kMsl;
}
#endif // TINT_BUILD_WGSL_WRITER
#endif // TINT_BUILD_MSL_WRITER
#if TINT_BUILD_HLSL_WRITER
if (ends_with(filename, ".hlsl")) {
return Format::kHlsl;
}
#endif // TINT_BUILD_HLSL_WRITER
return Format::kNone;
}

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

@ -390,6 +390,16 @@ std::string GetGlslStd450FuncName(uint32_t ext_opcode) {
return "mix";
case GLSLstd450Normalize:
return "normalize";
case GLSLstd450PackSnorm4x8:
return "pack4x8snorm";
case GLSLstd450PackUnorm4x8:
return "pack4x8unorm";
case GLSLstd450PackSnorm2x16:
return "pack2x16snorm";
case GLSLstd450PackUnorm2x16:
return "pack2x16unorm";
case GLSLstd450PackHalf2x16:
return "pack2x16float";
case GLSLstd450Pow:
return "pow";
case GLSLstd450FSign:
@ -437,11 +447,6 @@ std::string GetGlslStd450FuncName(uint32_t ext_opcode) {
case GLSLstd450Frexp:
case GLSLstd450FrexpStruct:
case GLSLstd450PackSnorm4x8:
case GLSLstd450PackUnorm4x8:
case GLSLstd450PackSnorm2x16:
case GLSLstd450PackUnorm2x16:
case GLSLstd450PackHalf2x16:
case GLSLstd450PackDouble2x32:
case GLSLstd450UnpackSnorm2x16:
case GLSLstd450UnpackUnorm2x16:

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

@ -57,6 +57,7 @@ std::string Preamble() {
OpName %v2f3 "v2f3"
OpName %v3f1 "v3f1"
OpName %v3f2 "v3f2"
OpName %v4f1 "v4f1"
%void = OpTypeVoid
%voidfn = OpTypeFunction %void
@ -79,6 +80,7 @@ std::string Preamble() {
%v2int = OpTypeVector %int 2
%v2float = OpTypeVector %float 2
%v3float = OpTypeVector %float 3
%v4float = OpTypeVector %float 4
%v2uint_10_20 = OpConstantComposite %v2uint %uint_10 %uint_20
%v2uint_20_10 = OpConstantComposite %v2uint %uint_20 %uint_10
@ -93,6 +95,8 @@ std::string Preamble() {
%v3float_50_60_70 = OpConstantComposite %v3float %float_50 %float_60 %float_70
%v3float_60_70_50 = OpConstantComposite %v3float %float_60 %float_70 %float_50
%v4float_50_50_50_50 = OpConstantComposite %v4float %float_50 %float_50 %float_50 %float_50
%100 = OpFunction %void None %voidfn
%entry = OpLabel
@ -123,6 +127,7 @@ std::string Preamble() {
%v3f1 = OpCopyObject %v3float %v3float_50_60_70
%v3f2 = OpCopyObject %v3float %v3float_60_70_50
%v4f1 = OpCopyObject %v4float %v4float_50_50_50_50
)";
}
@ -183,22 +188,24 @@ TEST_P(SpvParserTest_GlslStd450_Float_Floating, Scalar) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << assembly;
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__f32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{f1}
)
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
TEST_P(SpvParserTest_GlslStd450_Float_Floating, Vector) {
@ -212,22 +219,24 @@ TEST_P(SpvParserTest_GlslStd450_Float_Floating, Vector) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__f32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{v2f1}
)
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
TEST_P(SpvParserTest_GlslStd450_Float_FloatingFloating, Scalar) {
@ -241,15 +250,17 @@ TEST_P(SpvParserTest_GlslStd450_Float_FloatingFloating, Scalar) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << assembly;
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__f32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{f1}
Identifier[not set]{f2}
@ -257,7 +268,7 @@ TEST_P(SpvParserTest_GlslStd450_Float_FloatingFloating, Scalar) {
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
TEST_P(SpvParserTest_GlslStd450_Float_FloatingFloating, Vector) {
@ -271,15 +282,17 @@ TEST_P(SpvParserTest_GlslStd450_Float_FloatingFloating, Vector) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__f32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{v2f1}
Identifier[not set]{v2f2}
@ -287,7 +300,7 @@ TEST_P(SpvParserTest_GlslStd450_Float_FloatingFloating, Vector) {
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
TEST_P(SpvParserTest_GlslStd450_Floating_Floating, Scalar) {
@ -301,22 +314,24 @@ TEST_P(SpvParserTest_GlslStd450_Floating_Floating, Scalar) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << assembly;
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__f32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{f1}
)
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
TEST_P(SpvParserTest_GlslStd450_Floating_Floating, Vector) {
@ -330,22 +345,24 @@ TEST_P(SpvParserTest_GlslStd450_Floating_Floating, Vector) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__vec_2__f32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{v2f1}
)
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
TEST_P(SpvParserTest_GlslStd450_Floating_FloatingFloating, Scalar) {
@ -359,15 +376,17 @@ TEST_P(SpvParserTest_GlslStd450_Floating_FloatingFloating, Scalar) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << assembly;
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__f32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{f1}
Identifier[not set]{f2}
@ -375,7 +394,7 @@ TEST_P(SpvParserTest_GlslStd450_Floating_FloatingFloating, Scalar) {
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
TEST_P(SpvParserTest_GlslStd450_Floating_FloatingFloating, Vector) {
@ -389,15 +408,17 @@ TEST_P(SpvParserTest_GlslStd450_Floating_FloatingFloating, Vector) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << assembly;
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__vec_2__f32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{v2f1}
Identifier[not set]{v2f2}
@ -405,7 +426,7 @@ TEST_P(SpvParserTest_GlslStd450_Floating_FloatingFloating, Vector) {
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
TEST_P(SpvParserTest_GlslStd450_Floating_FloatingFloatingFloating, Scalar) {
@ -419,15 +440,17 @@ TEST_P(SpvParserTest_GlslStd450_Floating_FloatingFloatingFloating, Scalar) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__f32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{f1}
Identifier[not set]{f2}
@ -436,7 +459,7 @@ TEST_P(SpvParserTest_GlslStd450_Floating_FloatingFloatingFloating, Scalar) {
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
TEST_P(SpvParserTest_GlslStd450_Floating_FloatingFloatingFloating, Vector) {
@ -451,15 +474,17 @@ TEST_P(SpvParserTest_GlslStd450_Floating_FloatingFloatingFloating, Vector) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__vec_2__f32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{v2f1}
Identifier[not set]{v2f2}
@ -468,7 +493,7 @@ TEST_P(SpvParserTest_GlslStd450_Floating_FloatingFloatingFloating, Vector) {
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
TEST_P(SpvParserTest_GlslStd450_Floating_FloatingUinting, Scalar) {
@ -482,15 +507,17 @@ TEST_P(SpvParserTest_GlslStd450_Floating_FloatingUinting, Scalar) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__f32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{f1}
Identifier[not set]{u1}
@ -498,7 +525,7 @@ TEST_P(SpvParserTest_GlslStd450_Floating_FloatingUinting, Scalar) {
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
TEST_P(SpvParserTest_GlslStd450_Floating_FloatingUinting, Vector) {
@ -513,15 +540,17 @@ TEST_P(SpvParserTest_GlslStd450_Floating_FloatingUinting, Vector) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__vec_2__f32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{v2f1}
Identifier[not set]{v2u1}
@ -529,7 +558,7 @@ TEST_P(SpvParserTest_GlslStd450_Floating_FloatingUinting, Vector) {
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
TEST_P(SpvParserTest_GlslStd450_Floating_FloatingInting, Scalar) {
@ -543,15 +572,17 @@ TEST_P(SpvParserTest_GlslStd450_Floating_FloatingInting, Scalar) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__f32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{f1}
Identifier[not set]{i1}
@ -559,7 +590,7 @@ TEST_P(SpvParserTest_GlslStd450_Floating_FloatingInting, Scalar) {
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
TEST_P(SpvParserTest_GlslStd450_Floating_FloatingInting, Vector) {
@ -574,15 +605,17 @@ TEST_P(SpvParserTest_GlslStd450_Floating_FloatingInting, Vector) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__vec_2__f32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{v2f1}
Identifier[not set]{v2i1}
@ -590,7 +623,7 @@ TEST_P(SpvParserTest_GlslStd450_Floating_FloatingInting, Vector) {
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
TEST_P(SpvParserTest_GlslStd450_Float3_Float3Float3, Samples) {
@ -605,15 +638,17 @@ TEST_P(SpvParserTest_GlslStd450_Float3_Float3Float3, Samples) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__vec_3__f32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{v3f1}
Identifier[not set]{v3f2}
@ -621,7 +656,7 @@ TEST_P(SpvParserTest_GlslStd450_Float3_Float3Float3, Samples) {
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
INSTANTIATE_TEST_SUITE_P(Samples,
@ -709,22 +744,24 @@ TEST_P(SpvParserTest_GlslStd450_Inting_Inting, Scalar) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__i32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{i1}
)
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
TEST_P(SpvParserTest_GlslStd450_Inting_Inting, Vector) {
@ -739,22 +776,24 @@ TEST_P(SpvParserTest_GlslStd450_Inting_Inting, Vector) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__vec_2__i32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{v2i1}
)
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
TEST_P(SpvParserTest_GlslStd450_Inting_IntingInting, Scalar) {
@ -769,15 +808,17 @@ TEST_P(SpvParserTest_GlslStd450_Inting_IntingInting, Scalar) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__i32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{i1}
Identifier[not set]{i2}
@ -785,7 +826,7 @@ TEST_P(SpvParserTest_GlslStd450_Inting_IntingInting, Scalar) {
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
TEST_P(SpvParserTest_GlslStd450_Inting_IntingInting, Vector) {
@ -800,15 +841,17 @@ TEST_P(SpvParserTest_GlslStd450_Inting_IntingInting, Vector) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__vec_2__i32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{v2i1}
Identifier[not set]{v2i2}
@ -816,7 +859,7 @@ TEST_P(SpvParserTest_GlslStd450_Inting_IntingInting, Vector) {
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
TEST_P(SpvParserTest_GlslStd450_Inting_IntingIntingInting, Scalar) {
@ -831,15 +874,17 @@ TEST_P(SpvParserTest_GlslStd450_Inting_IntingIntingInting, Scalar) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__i32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{i1}
Identifier[not set]{i2}
@ -848,7 +893,7 @@ TEST_P(SpvParserTest_GlslStd450_Inting_IntingIntingInting, Scalar) {
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
TEST_P(SpvParserTest_GlslStd450_Inting_IntingIntingInting, Vector) {
@ -863,15 +908,17 @@ TEST_P(SpvParserTest_GlslStd450_Inting_IntingIntingInting, Vector) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__vec_2__i32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{v2i1}
Identifier[not set]{v2i2}
@ -880,7 +927,7 @@ TEST_P(SpvParserTest_GlslStd450_Inting_IntingIntingInting, Vector) {
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
INSTANTIATE_TEST_SUITE_P(Samples,
@ -907,15 +954,17 @@ TEST_P(SpvParserTest_GlslStd450_Uinting_UintingUinting, Scalar) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__u32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{u1}
Identifier[not set]{u2}
@ -923,7 +972,7 @@ TEST_P(SpvParserTest_GlslStd450_Uinting_UintingUinting, Scalar) {
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
TEST_P(SpvParserTest_GlslStd450_Uinting_UintingUinting, Vector) {
@ -938,15 +987,17 @@ TEST_P(SpvParserTest_GlslStd450_Uinting_UintingUinting, Vector) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__vec_2__u32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{v2u1}
Identifier[not set]{v2u2}
@ -954,7 +1005,7 @@ TEST_P(SpvParserTest_GlslStd450_Uinting_UintingUinting, Vector) {
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
TEST_P(SpvParserTest_GlslStd450_Uinting_UintingUintingUinting, Scalar) {
@ -968,15 +1019,17 @@ TEST_P(SpvParserTest_GlslStd450_Uinting_UintingUintingUinting, Scalar) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__u32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{u1}
Identifier[not set]{u2}
@ -985,7 +1038,7 @@ TEST_P(SpvParserTest_GlslStd450_Uinting_UintingUintingUinting, Scalar) {
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
TEST_P(SpvParserTest_GlslStd450_Uinting_UintingUintingUinting, Vector) {
@ -1000,15 +1053,17 @@ TEST_P(SpvParserTest_GlslStd450_Uinting_UintingUintingUinting, Vector) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
EXPECT_THAT(ToString(p->builder(), fe.ast_body()), HasSubstr(R"(
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__vec_2__u32
{
Call[not set]{
Identifier[not set]{)" + GetParam().wgsl_func +
R"(}
Identifier[not set]{)" +
GetParam().wgsl_func +
R"(}
(
Identifier[not set]{v2u1}
Identifier[not set]{v2u2}
@ -1017,7 +1072,7 @@ TEST_P(SpvParserTest_GlslStd450_Uinting_UintingUintingUinting, Vector) {
}
}
})"))
<< ToString(p->builder(), fe.ast_body());
<< body;
}
INSTANTIATE_TEST_SUITE_P(Samples,
@ -1043,7 +1098,7 @@ TEST_F(SpvParserTest, RectifyOperandsAndResult_SAbs) {
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
auto body = ToString(p->builder(), fe.ast_body());
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
@ -1436,6 +1491,62 @@ TEST_F(SpvParserTest, RectifyOperandsAndResult_UClamp) {
<< body;
}
struct DataPackingCase {
std::string opcode;
std::string wgsl_func;
uint32_t vec_size;
};
inline std::ostream& operator<<(std::ostream& out, DataPackingCase c) {
out << "DataPacking(" << c.opcode << ")";
return out;
}
using SpvParserTest_GlslStd450_DataPacking =
SpvParserTestBase<::testing::TestWithParam<DataPackingCase>>;
TEST_P(SpvParserTest_GlslStd450_DataPacking, Valid) {
auto param = GetParam();
const auto assembly = Preamble() + R"(
%1 = OpExtInst %uint %glsl )" +
param.opcode +
(param.vec_size == 2 ? " %v2f1" : " %v4f1") + R"(
OpReturn
OpFunctionEnd
)";
auto p = parser(test::Assemble(assembly));
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << assembly;
FunctionEmitter fe(p.get(), *spirv_function(p.get(), 100));
EXPECT_TRUE(fe.EmitBody()) << p->error();
const auto body = ToString(p->builder(), fe.ast_body());
EXPECT_THAT(body, HasSubstr(R"(
VariableConst{
x_1
none
__u32
{
Call[not set]{
Identifier[not set]{)" +
param.wgsl_func + R"(}
(
Identifier[not set]{v)" +
std::to_string(param.vec_size) + R"(f1}
)
}
}
})"))
<< body;
}
INSTANTIATE_TEST_SUITE_P(Samples,
SpvParserTest_GlslStd450_DataPacking,
::testing::ValuesIn(std::vector<DataPackingCase>{
{"PackSnorm4x8", "pack4x8snorm", 4},
{"PackUnorm4x8", "pack4x8unorm", 4},
{"PackSnorm2x16", "pack2x16snorm", 2},
{"PackUnorm2x16", "pack2x16unorm", 2},
{"PackHalf2x16", "pack2x16float", 2}}));
} // namespace
} // namespace spirv
} // namespace reader

16
src/semantic/intrinsic.cc
View File

@ -124,6 +124,16 @@ const char* str(Intrinsic i) {
return "modf";
case Intrinsic::kNormalize:
return "normalize";
case Intrinsic::kPack4x8Snorm:
return "pack4x8snorm";
case Intrinsic::kPack4x8Unorm:
return "pack4x8unorm";
case Intrinsic::kPack2x16Snorm:
return "pack2x16snorm";
case Intrinsic::kPack2x16Unorm:
return "pack2x16unorm";
case Intrinsic::kPack2x16Float:
return "pack2x16float";
case Intrinsic::kPow:
return "pow";
case Intrinsic::kReflect:
@ -215,6 +225,12 @@ bool IsImageQueryIntrinsic(Intrinsic i) {
i == Intrinsic::kTextureNumSamples;
}
bool IsDataPackingIntrinsic(Intrinsic i) {
return i == Intrinsic::kPack4x8Snorm || i == Intrinsic::kPack4x8Unorm ||
i == Intrinsic::kPack2x16Snorm || i == Intrinsic::kPack2x16Unorm ||
i == Intrinsic::kPack2x16Float;
}
} // namespace intrinsic
} // namespace semantic
} // namespace tint

10
src/semantic/intrinsic.h
View File

@ -70,6 +70,11 @@ enum class Intrinsic {
kMix,
kModf,
kNormalize,
kPack4x8Snorm,
kPack4x8Unorm,
kPack2x16Snorm,
kPack2x16Unorm,
kPack2x16Float,
kPow,
kReflect,
kReverseBits,
@ -133,6 +138,11 @@ bool IsTextureIntrinsic(Intrinsic i);
/// @returns true if the given `i` is a image query intrinsic
bool IsImageQueryIntrinsic(Intrinsic i);
/// Determines if the given `i` is a data packing intrinsic
/// @param i the intrinsic
/// @returns true if the given `i` is a data packing intrinsic
bool IsDataPackingIntrinsic(Intrinsic i);
/// @returns the name of the intrinsic function. The spelling, including case,
/// matches the name in the WGSL spec.
const char* str(Intrinsic i);

20
src/type_determiner.cc
View File

@ -517,6 +517,16 @@ constexpr const IntrinsicData kIntrinsicData[] = {
{semantic::Intrinsic::kMix, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kModf, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kNormalize, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kPack4x8Snorm, IntrinsicDataType::kUnsignedInteger, 1,
0},
{semantic::Intrinsic::kPack4x8Unorm, IntrinsicDataType::kUnsignedInteger, 1,
0},
{semantic::Intrinsic::kPack2x16Snorm, IntrinsicDataType::kUnsignedInteger,
1, 0},
{semantic::Intrinsic::kPack2x16Unorm, IntrinsicDataType::kUnsignedInteger,
1, 0},
{semantic::Intrinsic::kPack2x16Float, IntrinsicDataType::kUnsignedInteger,
1, 0},
{semantic::Intrinsic::kPow, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kReflect, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kReverseBits, IntrinsicDataType::kDependent, 0, 0},
@ -943,6 +953,16 @@ semantic::Intrinsic TypeDeterminer::MatchIntrinsic(const std::string& name) {
return semantic::Intrinsic::kModf;
} else if (name == "normalize") {
return semantic::Intrinsic::kNormalize;
} else if (name == "pack4x8snorm") {
return semantic::Intrinsic::kPack4x8Snorm;
} else if (name == "pack4x8unorm") {
return semantic::Intrinsic::kPack4x8Unorm;
} else if (name == "pack2x16snorm") {
return semantic::Intrinsic::kPack2x16Snorm;
} else if (name == "pack2x16unorm") {
return semantic::Intrinsic::kPack2x16Unorm;
} else if (name == "pack2x16float") {
return semantic::Intrinsic::kPack2x16Float;
} else if (name == "pow") {
return semantic::Intrinsic::kPow;
} else if (name == "reflect") {

23
src/type_determiner_test.cc
View File

@ -1623,6 +1623,29 @@ TEST_F(TypeDeterminerTest, MatchIntrinsicNoMatch) {
semantic::Intrinsic::kNone);
}
using ImportData_DataPackingTest = TypeDeterminerTestWithParam<IntrinsicData>;
TEST_P(ImportData_DataPackingTest, InferType) {
auto param = GetParam();
auto* ident = Expr(param.name);
auto* call = Call(ident);
WrapInFunction(call);
EXPECT_TRUE(td()->Determine()) << td()->error();
ASSERT_NE(TypeOf(call), nullptr);
EXPECT_TRUE(TypeOf(call)->Is<type::U32>());
}
INSTANTIATE_TEST_SUITE_P(
TypeDeterminerTest,
ImportData_DataPackingTest,
testing::Values(
IntrinsicData{"pack4x8snorm", semantic::Intrinsic::kPack4x8Snorm},
IntrinsicData{"pack4x8unorm", semantic::Intrinsic::kPack4x8Unorm},
IntrinsicData{"pack2x16snorm", semantic::Intrinsic::kPack2x16Snorm},
IntrinsicData{"pack2x16unorm", semantic::Intrinsic::kPack2x16Unorm},
IntrinsicData{"pack2x16float", semantic::Intrinsic::kPack2x16Float}));
using ImportData_SingleParamTest = TypeDeterminerTestWithParam<IntrinsicData>;
TEST_P(ImportData_SingleParamTest, Scalar) {
auto param = GetParam();

132
src/validator/validator_builtins_test.cc
View File

@ -1632,4 +1632,136 @@ INSTANTIATE_TEST_SUITE_P(ValidatorBuiltinsTest,
::testing::Values(std::make_tuple("all", 1),
std::make_tuple("any", 1)));
using DataPacking4x8 = ValidatorBuiltinsTestWithParams<std::string>;
TEST_P(DataPacking4x8, Float_Vec4) {
auto name = GetParam();
auto* builtin = Call(name, vec4<float>(1.0f, 1.0f, 1.0f, 1.0f));
WrapInFunction(builtin);
ValidatorImpl& v = Build();
EXPECT_TRUE(v.ValidateCallExpr(builtin)) << v.error();
}
TEST_P(DataPacking4x8, Float_Vec2) {
auto name = GetParam();
auto* builtin = Call(name, vec2<float>(1.0f, 1.0f));
WrapInFunction(builtin);
ValidatorImpl& v = Build();
EXPECT_FALSE(v.ValidateCallExpr(builtin));
EXPECT_EQ(v.error(),
"incorrect vector size for " + name + ". Requires 4 elements");
}
TEST_P(DataPacking4x8, Int_Vec4) {
auto name = GetParam();
auto* builtin = Call(name, vec4<int>(1, 1, 1, 1));
WrapInFunction(builtin);
ValidatorImpl& v = Build();
EXPECT_FALSE(v.ValidateCallExpr(builtin));
EXPECT_EQ(v.error(),
"incorrect type for " + name + ". Requires float vector value");
}
TEST_P(DataPacking4x8, Float_Scalar) {
auto name = GetParam();
auto* builtin = Call(name, Expr(1.0f));
WrapInFunction(builtin);
ValidatorImpl& v = Build();
EXPECT_FALSE(v.ValidateCallExpr(builtin));
EXPECT_EQ(v.error(),
"incorrect type for " + name + ". Requires float vector value");
}
TEST_P(DataPacking4x8, TooFewParams) {
auto name = GetParam();
auto* builtin = Call(name);
WrapInFunction(builtin);
ValidatorImpl& v = Build();
EXPECT_FALSE(v.ValidateCallExpr(builtin));
EXPECT_EQ(v.error(),
"incorrect number of parameters for " + name + " expected 1 got 0");
}
TEST_P(DataPacking4x8, TooManyParams) {
auto name = GetParam();
auto* builtin = Call(name, vec4<float>(1.0f, 1.0f, 1.0f, 1.0f),
vec4<float>(1.0f, 1.0f, 1.0f, 1.0f));
WrapInFunction(builtin);
ValidatorImpl& v = Build();
EXPECT_FALSE(v.ValidateCallExpr(builtin));
EXPECT_EQ(v.error(),
"incorrect number of parameters for " + name + " expected 1 got 2");
}
INSTANTIATE_TEST_SUITE_P(ValidatorBuiltinsTest,
DataPacking4x8,
::testing::Values("pack4x8snorm", "pack4x8unorm"));
using DataPacking2x16 = ValidatorBuiltinsTestWithParams<std::string>;
TEST_P(DataPacking2x16, Float_Vec4) {
auto name = GetParam();
auto* builtin = Call(name, vec4<float>(1.0f, 1.0f, 1.0f, 1.0f));
WrapInFunction(builtin);
ValidatorImpl& v = Build();
EXPECT_FALSE(v.ValidateCallExpr(builtin));
EXPECT_EQ(v.error(),
"incorrect vector size for " + name + ". Requires 2 elements");
}
TEST_P(DataPacking2x16, Float_Vec2) {
auto name = GetParam();
auto* builtin = Call(name, vec2<float>(1.0f, 1.0f));
WrapInFunction(builtin);
ValidatorImpl& v = Build();
EXPECT_TRUE(v.ValidateCallExpr(builtin)) << v.error();
}
TEST_P(DataPacking2x16, Int_Vec4) {
auto name = GetParam();
auto* builtin = Call(name, vec4<int>(1, 1, 1, 1));
WrapInFunction(builtin);
ValidatorImpl& v = Build();
EXPECT_FALSE(v.ValidateCallExpr(builtin));
EXPECT_EQ(v.error(),
"incorrect type for " + name + ". Requires float vector value");
}
TEST_P(DataPacking2x16, Float_Scalar) {
auto name = GetParam();
auto* builtin = Call(name, Expr(1.0f));
WrapInFunction(builtin);
ValidatorImpl& v = Build();
EXPECT_FALSE(v.ValidateCallExpr(builtin));
EXPECT_EQ(v.error(),
"incorrect type for " + name + ". Requires float vector value");
}
TEST_P(DataPacking2x16, TooFewParams) {
auto name = GetParam();
auto* builtin = Call(name);
WrapInFunction(builtin);
ValidatorImpl& v = Build();
EXPECT_FALSE(v.ValidateCallExpr(builtin));
EXPECT_EQ(v.error(),
"incorrect number of parameters for " + name + " expected 1 got 0");
}
TEST_P(DataPacking2x16, TooManyParams) {
auto name = GetParam();
auto* builtin = Call(name, vec4<float>(1.0f, 1.0f, 1.0f, 1.0f),
vec4<float>(1.0f, 1.0f, 1.0f, 1.0f));
WrapInFunction(builtin);
ValidatorImpl& v = Build();
EXPECT_FALSE(v.ValidateCallExpr(builtin));
EXPECT_EQ(v.error(),
"incorrect number of parameters for " + name + " expected 1 got 2");
}
INSTANTIATE_TEST_SUITE_P(ValidatorBuiltinsTest,
DataPacking2x16,
::testing::Values("pack2x16snorm",
"pack2x16unorm",
"pack2x16float"));
} // namespace tint

19
src/validator/validator_impl.cc
View File

@ -153,6 +153,16 @@ constexpr const IntrinsicData kIntrinsicData[] = {
true},
{semantic::Intrinsic::kNormalize, 1, IntrinsicDataType::kFloatVector, 0,
true},
{semantic::Intrinsic::kPack4x8Snorm, 1, IntrinsicDataType::kFloatVector, 4,
false},
{semantic::Intrinsic::kPack4x8Unorm, 1, IntrinsicDataType::kFloatVector, 4,
false},
{semantic::Intrinsic::kPack2x16Snorm, 1, IntrinsicDataType::kFloatVector, 2,
false},
{semantic::Intrinsic::kPack2x16Unorm, 1, IntrinsicDataType::kFloatVector, 2,
false},
{semantic::Intrinsic::kPack2x16Float, 1, IntrinsicDataType::kFloatVector, 2,
false},
{semantic::Intrinsic::kPow, 2, IntrinsicDataType::kFloatScalarOrVector, 0,
true},
{semantic::Intrinsic::kReflect, 2, IntrinsicDataType::kFloatScalarOrVector,
@ -845,6 +855,15 @@ bool ValidatorImpl::ValidateCallExpr(const ast::CallExpression* expr) {
}
}
if (semantic::intrinsic::IsDataPackingIntrinsic(data->intrinsic)) {
if (!program_->TypeOf(expr)->Is<type::U32>()) {
add_error(expr->source(),
"incorrect type for " + builtin +
". Result type must be an unsigned int scalar");
return false;
}
}
if (data->intrinsic == semantic::Intrinsic::kLength ||
data->intrinsic == semantic::Intrinsic::kDistance ||
data->intrinsic == semantic::Intrinsic::kDeterminant) {

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

@ -554,29 +554,30 @@ bool GeneratorImpl::EmitCall(std::ostream& pre,
} else if (sem->intrinsic() == semantic::Intrinsic::kIsNormal) {
error_ = "is_normal not supported in HLSL backend yet";
return false;
} else {
auto name = generate_builtin_name(sem);
if (name.empty()) {
} else if (semantic::intrinsic::IsDataPackingIntrinsic(sem->intrinsic())) {
return EmitDataPackingCall(pre, out, expr);
}
auto name = generate_builtin_name(sem);
if (name.empty()) {
return false;
}
make_indent(out);
out << name << "(";
bool first = true;
for (auto* param : params) {
if (!first) {
out << ", ";
}
first = false;
if (!EmitExpression(pre, out, param)) {
return false;
}
make_indent(out);
out << name << "(";
bool first = true;
for (auto* param : params) {
if (!first) {
out << ", ";
}
first = false;
if (!EmitExpression(pre, out, param)) {
return false;
}
}
out << ")";
}
out << ")";
return true;
}
@ -635,6 +636,65 @@ bool GeneratorImpl::EmitCall(std::ostream& pre,
return true;
}
bool GeneratorImpl::EmitDataPackingCall(std::ostream& pre,
std::ostream& out,
ast::CallExpression* expr) {
auto* ident = builder_.Sem().Get(expr)->As<semantic::IntrinsicCall>();
auto* param = expr->params()[0];
auto tmp_name = generate_name(kTempNamePrefix);
std::ostringstream expr_out;
if (!EmitExpression(pre, expr_out, param)) {
return false;
}
uint32_t dims = 2;
bool is_signed = false;
uint32_t scale = 65535;
if (ident->intrinsic() == semantic::Intrinsic::kPack4x8Snorm ||
ident->intrinsic() == semantic::Intrinsic::kPack4x8Unorm) {
dims = 4;
scale = 255;
}
if (ident->intrinsic() == semantic::Intrinsic::kPack4x8Snorm ||
ident->intrinsic() == semantic::Intrinsic::kPack2x16Snorm) {
is_signed = true;
scale = (scale - 1) / 2;
}
switch (ident->intrinsic()) {
case semantic::Intrinsic::kPack4x8Snorm:
case semantic::Intrinsic::kPack4x8Unorm:
case semantic::Intrinsic::kPack2x16Snorm:
case semantic::Intrinsic::kPack2x16Unorm:
pre << (is_signed ? "" : "u") << "int" << dims << " " << tmp_name << " = "
<< (is_signed ? "" : "u") << "int" << dims << "(round(clamp("
<< expr_out.str() << ", " << (is_signed ? "-1.0" : "0.0")
<< ", 1.0) * " << scale << ".0))";
if (is_signed) {
pre << " & " << (dims == 4 ? "0xff" : "0xffff");
}
pre << ";\n";
if (is_signed) {
out << "asuint";
}
out << "(";
out << tmp_name << ".x | " << tmp_name << ".y << " << (32 / dims);
if (dims == 4) {
out << " | " << tmp_name << ".z << 16 | " << tmp_name << ".w << 24";
}
out << ")";
break;
case semantic::Intrinsic::kPack2x16Float:
pre << "uint2 " << tmp_name << " = f32tof16(" << expr_out.str() << ");\n";
out << "(" << tmp_name << ".x | " << tmp_name << ".y << 16)";
break;
default:
error_ = "Internal error: unhandled data packing intrinsic";
return false;
}
return true;
}
bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
std::ostream& out,
ast::CallExpression* expr,
@ -1503,11 +1563,12 @@ bool GeneratorImpl::EmitEntryPointData(
return false;
}
if (ac->IsReadWrite()) {
if (!ac->IsReadOnly()) {
out << "RW";
}
out << "ByteAddressBuffer " << builder_.Symbols().NameFor(decl->symbol())
<< " : register(u" << binding->value() << ");" << std::endl;
<< " : register(" << (ac->IsReadOnly() ? "t" : "u") << binding->value()
<< ");" << std::endl;
emitted_storagebuffer = true;
}
if (emitted_storagebuffer) {

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

@ -159,6 +159,14 @@ class GeneratorImpl {
std::ostream& out,
ast::CallExpression* expr,
const semantic::TextureIntrinsicCall* sem);
/// Handles generating a call to data packing intrinsic
/// @param pre the preamble of the expression stream
/// @param out the output of the expression stream
/// @param expr the call expression
/// @returns true if the call expression is emitted
bool EmitDataPackingCall(std::ostream& pre,
std::ostream& out,
ast::CallExpression* expr);
/// Handles a case statement
/// @param out the output stream
/// @param stmt the statement

41
src/writer/hlsl/generator_impl_function_test.cc
View File

@ -407,7 +407,7 @@ TEST_F(HlslGeneratorImplTest_Function,
GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate(out)) << gen.error();
EXPECT_EQ(result(), R"(ByteAddressBuffer coord : register(u0);
EXPECT_EQ(result(), R"(ByteAddressBuffer coord : register(t0);
void frag_main() {
float v = asfloat(coord.Load(4));
@ -417,6 +417,45 @@ void frag_main() {
)");
}
TEST_F(HlslGeneratorImplTest_Function,
Emit_FunctionDecoration_EntryPoint_With_WO_StorageBuffer_Store) {
auto* str = create<ast::Struct>(
ast::StructMemberList{Member("a", ty.i32(), {MemberOffset(0)}),
Member("b", ty.f32(), {MemberOffset(4)})},
ast::StructDecorationList{});
auto* s = ty.struct_("Data", str);
type::AccessControl ac(ast::AccessControl::kWriteOnly, s);
Global("coord", ast::StorageClass::kStorage, &ac, nullptr,
ast::VariableDecorationList{
create<ast::BindingDecoration>(0),
create<ast::GroupDecoration>(1),
});
Func("frag_main", ast::VariableList{}, ty.void_(),
ast::StatementList{
create<ast::AssignmentStatement>(MemberAccessor("coord", "b"),
Expr(2.0f)),
create<ast::ReturnStatement>(),
},
ast::FunctionDecorationList{
create<ast::StageDecoration>(ast::PipelineStage::kFragment),
});
GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate(out)) << gen.error();
EXPECT_EQ(result(), R"(RWByteAddressBuffer coord : register(u0);
void frag_main() {
coord.Store(4, asuint(2.0f));
return;
}
)");
}
TEST_F(HlslGeneratorImplTest_Function,
Emit_FunctionDecoration_EntryPoint_With_StorageBuffer_Store) {
auto* str = create<ast::Struct>(

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

@ -14,6 +14,7 @@
#include <sstream>
#include "gmock/gmock.h"
#include "src/ast/call_expression.h"
#include "src/ast/identifier_expression.h"
#include "src/program.h"
@ -28,6 +29,8 @@ namespace writer {
namespace hlsl {
namespace {
using ::testing::HasSubstr;
using HlslGeneratorImplTest_Intrinsic = TestHelper;
enum class ParamType {
@ -279,6 +282,72 @@ TEST_F(HlslGeneratorImplTest_Intrinsic, Intrinsic_Call) {
EXPECT_EQ(result(), " dot(param1, param2)");
}
TEST_F(HlslGeneratorImplTest_Intrinsic, Pack4x8Snorm) {
auto* call = Call("pack4x8snorm", "p1");
Global("p1", ast::StorageClass::kPrivate, ty.vec4<f32>());
WrapInFunction(call);
GeneratorImpl& gen = Build();
gen.increment_indent();
ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error();
EXPECT_THAT(pre_result(), HasSubstr("int4 _tint_tmp = int4(round(clamp(p1, "
"-1.0, 1.0) * 127.0)) & 0xff;"));
EXPECT_THAT(result(), HasSubstr("asuint(_tint_tmp.x | _tint_tmp.y << 8 | "
"_tint_tmp.z << 16 | _tint_tmp.w << 24)"));
}
TEST_F(HlslGeneratorImplTest_Intrinsic, Pack4x8Unorm) {
auto* call = Call("pack4x8unorm", "p1");
Global("p1", ast::StorageClass::kPrivate, ty.vec4<f32>());
WrapInFunction(call);
GeneratorImpl& gen = Build();
gen.increment_indent();
ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error();
EXPECT_THAT(pre_result(), HasSubstr("uint4 _tint_tmp = uint4(round(clamp(p1, "
"0.0, 1.0) * 255.0));"));
EXPECT_THAT(result(), HasSubstr("(_tint_tmp.x | _tint_tmp.y << 8 | "
"_tint_tmp.z << 16 | _tint_tmp.w << 24)"));
}
TEST_F(HlslGeneratorImplTest_Intrinsic, Pack2x16Snorm) {
auto* call = Call("pack2x16snorm", "p1");
Global("p1", ast::StorageClass::kPrivate, ty.vec4<f32>());
WrapInFunction(call);
GeneratorImpl& gen = Build();
gen.increment_indent();
ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error();
EXPECT_THAT(pre_result(), HasSubstr("int2 _tint_tmp = int2(round(clamp(p1, "
"-1.0, 1.0) * 32767.0)) & 0xffff;"));
EXPECT_THAT(result(), HasSubstr("asuint(_tint_tmp.x | _tint_tmp.y << 16)"));
}
TEST_F(HlslGeneratorImplTest_Intrinsic, Pack2x16Unorm) {
auto* call = Call("pack2x16unorm", "p1");
Global("p1", ast::StorageClass::kPrivate, ty.vec4<f32>());
WrapInFunction(call);
GeneratorImpl& gen = Build();
gen.increment_indent();
ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error();
EXPECT_THAT(pre_result(), HasSubstr("uint2 _tint_tmp = uint2(round(clamp(p1, "
"0.0, 1.0) * 65535.0));"));
EXPECT_THAT(result(), HasSubstr("(_tint_tmp.x | _tint_tmp.y << 16)"));
}
TEST_F(HlslGeneratorImplTest_Intrinsic, Pack2x16float) {
auto* call = Call("pack2x16float", "p1");
Global("p1", ast::StorageClass::kPrivate, ty.vec4<f32>());
WrapInFunction(call);
GeneratorImpl& gen = Build();
gen.increment_indent();
ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error();
EXPECT_THAT(pre_result(), HasSubstr("uint2 _tint_tmp = f32tof16(p1);"));
EXPECT_THAT(result(), HasSubstr("(_tint_tmp.x | _tint_tmp.y << 16)"));
}
} // namespace
} // namespace hlsl
} // namespace writer

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

@ -449,6 +449,15 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) {
return EmitTextureCall(expr, sem);
}
if (auto* sem = call_sem->As<semantic::IntrinsicCall>()) {
if (sem->intrinsic() == semantic::Intrinsic::kPack2x16Float) {
make_indent();
out_ << "as_type<uint>(half2(";
if (!EmitExpression(expr->params()[0])) {
return false;
}
out_ << "))";
return true;
}
auto name = generate_builtin_name(sem);
if (name.empty()) {
return false;
@ -864,6 +873,18 @@ std::string GeneratorImpl::generate_builtin_name(
case semantic::Intrinsic::kFaceForward:
out += "faceforward";
break;
case semantic::Intrinsic::kPack4x8Snorm:
out += "pack_float_to_snorm4x8";
break;
case semantic::Intrinsic::kPack4x8Unorm:
out += "pack_float_to_unorm4x8";
break;
case semantic::Intrinsic::kPack2x16Snorm:
out += "pack_float_to_snorm2x16";
break;
case semantic::Intrinsic::kPack2x16Unorm:
out += "pack_float_to_unorm2x16";
break;
case semantic::Intrinsic::kReverseBits:
out += "reverse_bits";
break;

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

@ -96,6 +96,8 @@ ast::CallExpression* GenerateCall(semantic::Intrinsic intrinsic,
case semantic::Intrinsic::kLog:
case semantic::Intrinsic::kLog2:
case semantic::Intrinsic::kNormalize:
case semantic::Intrinsic::kPack4x8Snorm:
case semantic::Intrinsic::kPack4x8Unorm:
case semantic::Intrinsic::kReflect:
case semantic::Intrinsic::kRound:
case semantic::Intrinsic::kSin:
@ -147,6 +149,9 @@ ast::CallExpression* GenerateCall(semantic::Intrinsic intrinsic,
return builder->Call(str.str(), "f1", "f2", "b1");
case semantic::Intrinsic::kDeterminant:
return builder->Call(str.str(), "m1");
case semantic::Intrinsic::kPack2x16Snorm:
case semantic::Intrinsic::kPack2x16Unorm:
return builder->Call(str.str(), "f4");
default:
break;
}
@ -164,6 +169,7 @@ TEST_P(MslIntrinsicTest, Emit) {
Global("f1", ast::StorageClass::kFunction, ty.vec2<float>());
Global("f2", ast::StorageClass::kFunction, ty.vec2<float>());
Global("f3", ast::StorageClass::kFunction, ty.vec2<float>());
Global("f4", ast::StorageClass::kFunction, ty.vec2<float>());
Global("u1", ast::StorageClass::kFunction, ty.vec2<unsigned int>());
Global("u2", ast::StorageClass::kFunction, ty.vec2<unsigned int>());
Global("u3", ast::StorageClass::kFunction, ty.vec2<unsigned int>());
@ -269,6 +275,14 @@ INSTANTIATE_TEST_SUITE_P(
IntrinsicData{semantic::Intrinsic::kMin, ParamType::kU32, "metal::min"},
IntrinsicData{semantic::Intrinsic::kNormalize, ParamType::kF32,
"metal::normalize"},
IntrinsicData{semantic::Intrinsic::kPack4x8Snorm, ParamType::kF32,
"metal::pack_float_to_snorm4x8"},
IntrinsicData{semantic::Intrinsic::kPack4x8Unorm, ParamType::kF32,
"metal::pack_float_to_unorm4x8"},
IntrinsicData{semantic::Intrinsic::kPack2x16Snorm, ParamType::kF32,
"metal::pack_float_to_snorm2x16"},
IntrinsicData{semantic::Intrinsic::kPack2x16Unorm, ParamType::kF32,
"metal::pack_float_to_unorm2x16"},
IntrinsicData{semantic::Intrinsic::kPow, ParamType::kF32, "metal::pow"},
IntrinsicData{semantic::Intrinsic::kReflect, ParamType::kF32,
"metal::reflect"},
@ -309,6 +323,18 @@ TEST_F(MslGeneratorImplTest, Intrinsic_Call) {
EXPECT_EQ(gen.result(), " metal::dot(param1, param2)");
}
TEST_F(MslGeneratorImplTest, Pack2x16Float) {
auto* call = Call("pack2x16float", "p1");
Global("p1", ast::StorageClass::kFunction, ty.vec2<f32>());
WrapInFunction(call);
GeneratorImpl& gen = Build();
gen.increment_indent();
ASSERT_TRUE(gen.EmitExpression(call)) << gen.error();
EXPECT_EQ(gen.result(), " as_type<uint>(half2(p1))");
}
} // namespace
} // namespace msl
} // namespace writer

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

@ -249,6 +249,16 @@ uint32_t intrinsic_to_glsl_method(type::Type* type,
return GLSLstd450Modf;
case semantic::Intrinsic::kNormalize:
return GLSLstd450Normalize;
case semantic::Intrinsic::kPack4x8Snorm:
return GLSLstd450PackSnorm4x8;
case semantic::Intrinsic::kPack4x8Unorm:
return GLSLstd450PackUnorm4x8;
case semantic::Intrinsic::kPack2x16Snorm:
return GLSLstd450PackSnorm2x16;
case semantic::Intrinsic::kPack2x16Unorm:
return GLSLstd450PackUnorm2x16;
case semantic::Intrinsic::kPack2x16Float:
return GLSLstd450PackHalf2x16;
case semantic::Intrinsic::kPow:
return GLSLstd450Pow;
case semantic::Intrinsic::kReflect: