Disallow taking the address of a vector component

Update or remove tests that try to do this.

Fixed: tint:491
Change-Id: I1f351a4abf68ae9bc6b100885fb1bcea08b31211
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/68242
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>

docs/origin-trial-changes.md

@@ -1,5 +1,11 @@
 # Tint changes during Origin Trial
 
+## Changes for M98
+
+### Breaking Changes
+
+* Taking the address of a vector component is no longer allowed.
+
 ## Changes for M97
 
 ### Breaking Changes

src/resolver/ptr_ref_validation_test.cc

@@ -66,6 +66,34 @@ TEST_F(ResolverPtrRefValidationTest, AddressOfHandle) {
             "storage class");
 }
 
+TEST_F(ResolverPtrRefValidationTest, AddressOfVectorComponent_MemberAccessor) {
+  // var v : vec4<i32>;
+  // &v.y
+  auto* v = Var("v", ty.vec4<i32>());
+  auto* expr = AddressOf(MemberAccessor(Source{{12, 34}}, "v", "y"));
+
+  WrapInFunction(v, expr);
+
+  EXPECT_FALSE(r()->Resolve());
+
+  EXPECT_EQ(r()->error(),
+            "12:34 error: cannot take the address of a vector component");
+}
+
+TEST_F(ResolverPtrRefValidationTest, AddressOfVectorComponent_ArrayAccessor) {
+  // var v : vec4<i32>;
+  // &v[2]
+  auto* v = Var("v", ty.vec4<i32>());
+  auto* expr = AddressOf(IndexAccessor(Source{{12, 34}}, "v", 2));
+
+  WrapInFunction(v, expr);
+
+  EXPECT_FALSE(r()->Resolve());
+
+  EXPECT_EQ(r()->error(),
+            "12:34 error: cannot take the address of a vector component");
+}
+
 TEST_F(ResolverPtrRefValidationTest, IndirectOfAddressOfHandle) {
   // [[group(0), binding(0)]] var t: texture_3d<f32>;
   // *&t

src/resolver/resolver.cc

@@ -3428,6 +3428,17 @@ bool Resolver::UnaryOp(const ast::UnaryOpExpression* unary) {
               unary->expr->source);
           return false;
         }
+
+        auto* array = unary->expr->As<ast::ArrayAccessorExpression>();
+        auto* member = unary->expr->As<ast::MemberAccessorExpression>();
+        if ((array && TypeOf(array->array)->UnwrapRef()->Is<sem::Vector>()) ||
+            (member &&
+             TypeOf(member->structure)->UnwrapRef()->Is<sem::Vector>())) {
+          AddError("cannot take the address of a vector component",
+                   unary->expr->source);
+          return false;
+        }
+
         type = builder_->create<sem::Pointer>(
             ref->StoreType(), ref->StorageClass(), ref->Access());
       } else {

src/transform/inline_pointer_lets_test.cc

@@ -59,18 +59,18 @@ fn f() {
 TEST_F(InlinePointerLetsTest, ComplexChain) {
   auto* src = R"(
 fn f() {
-  var m : mat4x4<f32>;
+  var a : array<mat4x4<f32>, 4>;
-  let mp : ptr<function, mat4x4<f32>> = &m;
+  let ap : ptr<function, array<mat4x4<f32>, 4>> = &a;
+  let mp : ptr<function, mat4x4<f32>> = &(*ap)[3];
   let vp : ptr<function, vec4<f32>> = &(*mp)[2];
-  let fp : ptr<function, f32> = &(*vp)[1];
+  let v : vec4<f32> = *vp;
-  let f : f32 = *fp;
 }
 )";
 
   auto* expect = R"(
 fn f() {
-  var m : mat4x4<f32>;
+  var a : array<mat4x4<f32>, 4>;
-  let f : f32 = *(&((*(&((*(&(m)))[2])))[1]));
+  let v : vec4<f32> = *(&((*(&((*(&(a)))[3])))[2]));
 }
 )";
 
@@ -94,15 +94,6 @@ fn arr() {
   *p = 4;
 }
 
-fn vector() {
-  var v : vec3<f32>;
-  var i : i32 = 0;
-  var j : i32 = 0;
-  let p : ptr<function, f32> = &v[i + j];
-  i = 2;
-  *p = 4.0;
-}
-
 fn matrix() {
   var m : mat3x3<f32>;
   var i : i32 = 0;
@@ -127,22 +118,13 @@ fn arr() {
   *(&(a[p_save].i)) = 4;
 }
 
-fn vector() {
-  var v : vec3<f32>;
-  var i : i32 = 0;
-  var j : i32 = 0;
-  let p_save_1 = (i + j);
-  i = 2;
-  *(&(v[p_save_1])) = 4.0;
-}
-
 fn matrix() {
   var m : mat3x3<f32>;
   var i : i32 = 0;
   var j : i32 = 0;
-  let p_save_2 = (i + j);
+  let p_save_1 = (i + j);
   i = 2;
-  *(&(m[p_save_2])) = vec3<f32>(4.0, 5.0, 6.0);
+  *(&(m[p_save_1])) = vec3<f32>(4.0, 5.0, 6.0);
 }
 )";
 

src/writer/glsl/generator_impl_sanitizer_test.cc

@@ -302,29 +302,31 @@ void main() {
 }
 
 TEST_F(GlslSanitizerTest, InlinePtrLetsComplexChain) {
-  // var m : mat4x4<f32>;
+  // var a : array<mat4x4<f32>, 4>;
-  // let mp : ptr<function, mat4x4<f32>> = &m;
+  // let ap : ptr<function, array<mat4x4<f32>, 4>> = &a;
+  // let mp : ptr<function, mat4x4<f32>> = &(*ap)[3];
   // let vp : ptr<function, vec4<f32>> = &(*mp)[2];
-  // let fp : ptr<function, f32> = &(*vp)[1];
+  // let v : vec4<f32> = *vp;
-  // let f : f32 = *fp;
+  auto* a = Var("a", ty.array(ty.mat4x4<f32>(), 4));
-  auto* m = Var("m", ty.mat4x4<f32>());
+  auto* ap = Const(
+      "ap",
+      ty.pointer(ty.array(ty.mat4x4<f32>(), 4), ast::StorageClass::kFunction),
+      AddressOf(a));
   auto* mp =
       Const("mp", ty.pointer(ty.mat4x4<f32>(), ast::StorageClass::kFunction),
-            AddressOf(m));
+            AddressOf(IndexAccessor(Deref(ap), 3)));
   auto* vp =
       Const("vp", ty.pointer(ty.vec4<f32>(), ast::StorageClass::kFunction),
             AddressOf(IndexAccessor(Deref(mp), 2)));
-  auto* fp = Const("fp", ty.pointer<f32>(ast::StorageClass::kFunction),
+  auto* v = Var("v", ty.vec4<f32>(), ast::StorageClass::kNone, Deref(vp));
-                   AddressOf(IndexAccessor(Deref(vp), 1)));
-  auto* f = Var("f", ty.f32(), ast::StorageClass::kNone, Deref(fp));
 
   Func("main", ast::VariableList{}, ty.void_(),
        {
-           Decl(m),
+           Decl(a),
+           Decl(ap),
            Decl(mp),
            Decl(vp),
-           Decl(fp),
+           Decl(v),
-           Decl(f),
        },
        {
            Stage(ast::PipelineStage::kFragment),
@@ -339,8 +341,8 @@ TEST_F(GlslSanitizerTest, InlinePtrLetsComplexChain) {
 precision mediump float;
 
 void tint_symbol() {
-  mat4 m = mat4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+  mat4 a[4] = mat4[4](mat4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f), mat4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f), mat4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f), mat4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
-  float f = m[2][1];
+  vec4 v = a[3][2];
   return;
 }
 void main() {

src/writer/hlsl/generator_impl_sanitizer_test.cc

@@ -244,29 +244,31 @@ TEST_F(HlslSanitizerTest, InlinePtrLetsBasic) {
 }
 
 TEST_F(HlslSanitizerTest, InlinePtrLetsComplexChain) {
-  // var m : mat4x4<f32>;
+  // var a : array<mat4x4<f32>, 4>;
-  // let mp : ptr<function, mat4x4<f32>> = &m;
+  // let ap : ptr<function, array<mat4x4<f32>, 4>> = &a;
+  // let mp : ptr<function, mat4x4<f32>> = &(*ap)[3];
   // let vp : ptr<function, vec4<f32>> = &(*mp)[2];
-  // let fp : ptr<function, f32> = &(*vp)[1];
+  // let v : vec4<f32> = *vp;
-  // let f : f32 = *fp;
+  auto* a = Var("a", ty.array(ty.mat4x4<f32>(), 4));
-  auto* m = Var("m", ty.mat4x4<f32>());
+  auto* ap = Const(
+      "ap",
+      ty.pointer(ty.array(ty.mat4x4<f32>(), 4), ast::StorageClass::kFunction),
+      AddressOf(a));
   auto* mp =
       Const("mp", ty.pointer(ty.mat4x4<f32>(), ast::StorageClass::kFunction),
-            AddressOf(m));
+            AddressOf(IndexAccessor(Deref(ap), 3)));
   auto* vp =
       Const("vp", ty.pointer(ty.vec4<f32>(), ast::StorageClass::kFunction),
             AddressOf(IndexAccessor(Deref(mp), 2)));
-  auto* fp = Const("fp", ty.pointer<f32>(ast::StorageClass::kFunction),
+  auto* v = Var("v", ty.vec4<f32>(), ast::StorageClass::kNone, Deref(vp));
-                   AddressOf(IndexAccessor(Deref(vp), 1)));
-  auto* f = Var("f", ty.f32(), ast::StorageClass::kNone, Deref(fp));
 
   Func("main", ast::VariableList{}, ty.void_(),
        {
-           Decl(m),
+           Decl(a),
+           Decl(ap),
            Decl(mp),
            Decl(vp),
-           Decl(fp),
+           Decl(v),
-           Decl(f),
        },
        {
            Stage(ast::PipelineStage::kFragment),
@@ -278,8 +280,8 @@ TEST_F(HlslSanitizerTest, InlinePtrLetsComplexChain) {
 
   auto got = gen.result();
   auto* expect = R"(void main() {
-  float4x4 m = float4x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+  float4x4 a[4] = (float4x4[4])0;
-  float f = m[2][1];
+  float4 v = a[3][2];
   return;
 }
 )";

test/bug/tint/453.wgsl

@@ -6,8 +6,7 @@ fn main() {
   var srcValue : vec4<u32>;
   let x_22 : vec4<u32> = textureLoad(Src, vec2<i32>(0, 0), 0);
   srcValue = x_22;
-  let x_23 : ptr<function, u32> = &srcValue.x;
+  let x_24 : u32 = srcValue.x;
-  let x_24 : u32 = *x_23;
   let x_25 : u32 = (x_24 + 1u);
   let x_27 : vec4<u32> = srcValue;
   textureStore(Dst, vec2<i32>(0, 0), x_27.xxxx);

test/bug/tint/453.wgsl.expected.wgsl

@@ -7,8 +7,7 @@ fn main() {
   var srcValue : vec4<u32>;
   let x_22 : vec4<u32> = textureLoad(Src, vec2<i32>(0, 0), 0);
   srcValue = x_22;
-  let x_23 : ptr<function, u32> = &(srcValue.x);
+  let x_24 : u32 = srcValue.x;
-  let x_24 : u32 = *(x_23);
   let x_25 : u32 = (x_24 + 1u);
   let x_27 : vec4<u32> = srcValue;
   textureStore(Dst, vec2<i32>(0, 0), x_27.xxxx);

test/ptr_ref/access/vector.spvasm

@@ -1,33 +0,0 @@
-; SPIR-V
-; Version: 1.3
-; Generator: Google Tint Compiler; 0
-; Bound: 21
-; Schema: 0
-               OpCapability Shader
-               OpMemoryModel Logical GLSL450
-               OpEntryPoint GLCompute %main "main"
-               OpExecutionMode %main LocalSize 1 1 1
-               OpName %main "main"
-               OpName %v "v"
-       %void = OpTypeVoid
-          %1 = OpTypeFunction %void
-      %float = OpTypeFloat 32
-    %v3float = OpTypeVector %float 3
-    %float_1 = OpConstant %float 1
-    %float_2 = OpConstant %float 2
-    %float_3 = OpConstant %float 3
-         %10 = OpConstantComposite %v3float %float_1 %float_2 %float_3
-%_ptr_Function_v3float = OpTypePointer Function %v3float
-         %13 = OpConstantNull %v3float
-       %uint = OpTypeInt 32 0
-     %uint_1 = OpConstant %uint 1
-%_ptr_Function_float = OpTypePointer Function %float
-    %float_5 = OpConstant %float 5
-       %main = OpFunction %void None %1
-          %4 = OpLabel
-          %v = OpVariable %_ptr_Function_v3float Function %13
-               OpStore %v %10
-         %18 = OpAccessChain %_ptr_Function_float %v %uint_1
-               OpStore %18 %float_5
-               OpReturn
-               OpFunctionEnd

test/ptr_ref/access/vector.spvasm.expected.hlsl

@@ -1,12 +0,0 @@
-void main_1() {
-  float3 v = float3(0.0f, 0.0f, 0.0f);
-  v = float3(1.0f, 2.0f, 3.0f);
-  v.y = 5.0f;
-  return;
-}
-
-[numthreads(1, 1, 1)]
-void main() {
-  main_1();
-  return;
-}

test/ptr_ref/access/vector.spvasm.expected.msl

@@ -1,15 +0,0 @@
-#include <metal_stdlib>
-
-using namespace metal;
-void main_1() {
-  float3 v = float3(0.0f, 0.0f, 0.0f);
-  v = float3(1.0f, 2.0f, 3.0f);
-  v[1] = 5.0f;
-  return;
-}
-
-kernel void tint_symbol() {
-  main_1();
-  return;
-}
-

test/ptr_ref/access/vector.spvasm.expected.spvasm

@@ -1,42 +0,0 @@
-; SPIR-V
-; Version: 1.3
-; Generator: Google Tint Compiler; 0
-; Bound: 24
-; Schema: 0
-               OpCapability Shader
-               OpMemoryModel Logical GLSL450
-               OpEntryPoint GLCompute %main "main"
-               OpExecutionMode %main LocalSize 1 1 1
-               OpName %main_1 "main_1"
-               OpName %v "v"
-               OpName %main "main"
-       %void = OpTypeVoid
-          %1 = OpTypeFunction %void
-      %float = OpTypeFloat 32
-    %v3float = OpTypeVector %float 3
-    %float_0 = OpConstant %float 0
-          %8 = OpConstantComposite %v3float %float_0 %float_0 %float_0
-%_ptr_Function_v3float = OpTypePointer Function %v3float
-         %11 = OpConstantNull %v3float
-    %float_1 = OpConstant %float 1
-    %float_2 = OpConstant %float 2
-    %float_3 = OpConstant %float 3
-         %15 = OpConstantComposite %v3float %float_1 %float_2 %float_3
-       %uint = OpTypeInt 32 0
-     %uint_1 = OpConstant %uint 1
-%_ptr_Function_float = OpTypePointer Function %float
-    %float_5 = OpConstant %float 5
-     %main_1 = OpFunction %void None %1
-          %4 = OpLabel
-          %v = OpVariable %_ptr_Function_v3float Function %11
-               OpStore %v %8
-               OpStore %v %15
-         %19 = OpAccessChain %_ptr_Function_float %v %uint_1
-               OpStore %19 %float_5
-               OpReturn
-               OpFunctionEnd
-       %main = OpFunction %void None %1
-         %22 = OpLabel
-         %23 = OpFunctionCall %void %main_1
-               OpReturn
-               OpFunctionEnd

test/ptr_ref/access/vector.spvasm.expected.wgsl

@@ -1,11 +0,0 @@
-fn main_1() {
-  var v : vec3<f32> = vec3<f32>(0.0, 0.0, 0.0);
-  v = vec3<f32>(1.0, 2.0, 3.0);
-  v.y = 5.0;
-  return;
-}
-
-[[stage(compute), workgroup_size(1, 1, 1)]]
-fn main() {
-  main_1();
-}

test/ptr_ref/access/vector.wgsl

@@ -1,6 +0,0 @@
-[[stage(compute), workgroup_size(1)]]
-fn main() {
-  var v : vec3<f32> = vec3<f32>(1., 2., 3.);
-  let f : ptr<function, f32> = &v.y;
-  *f = 5.0;
-}

test/ptr_ref/access/vector.wgsl.expected.hlsl

@@ -1,6 +0,0 @@
-[numthreads(1, 1, 1)]
-void main() {
-  float3 v = float3(1.0f, 2.0f, 3.0f);
-  v.y = 5.0f;
-  return;
-}

test/ptr_ref/access/vector.wgsl.expected.msl

@@ -1,9 +0,0 @@
-#include <metal_stdlib>
-
-using namespace metal;
-kernel void tint_symbol() {
-  float3 v = float3(1.0f, 2.0f, 3.0f);
-  v[1] = 5.0f;
-  return;
-}
-

test/ptr_ref/access/vector.wgsl.expected.spvasm

@@ -1,33 +0,0 @@
-; SPIR-V
-; Version: 1.3
-; Generator: Google Tint Compiler; 0
-; Bound: 19
-; Schema: 0
-               OpCapability Shader
-               OpMemoryModel Logical GLSL450
-               OpEntryPoint GLCompute %main "main"
-               OpExecutionMode %main LocalSize 1 1 1
-               OpName %main "main"
-               OpName %v "v"
-       %void = OpTypeVoid
-          %1 = OpTypeFunction %void
-      %float = OpTypeFloat 32
-    %v3float = OpTypeVector %float 3
-    %float_1 = OpConstant %float 1
-    %float_2 = OpConstant %float 2
-    %float_3 = OpConstant %float 3
-         %10 = OpConstantComposite %v3float %float_1 %float_2 %float_3
-%_ptr_Function_v3float = OpTypePointer Function %v3float
-         %13 = OpConstantNull %v3float
-       %uint = OpTypeInt 32 0
-     %uint_1 = OpConstant %uint 1
-%_ptr_Function_float = OpTypePointer Function %float
-    %float_5 = OpConstant %float 5
-       %main = OpFunction %void None %1
-          %4 = OpLabel
-          %v = OpVariable %_ptr_Function_v3float Function %13
-               OpStore %v %10
-         %17 = OpAccessChain %_ptr_Function_float %v %uint_1
-               OpStore %17 %float_5
-               OpReturn
-               OpFunctionEnd

test/ptr_ref/access/vector.wgsl.expected.wgsl

@@ -1,6 +0,0 @@
-[[stage(compute), workgroup_size(1)]]
-fn main() {
-  var v : vec3<f32> = vec3<f32>(1.0, 2.0, 3.0);
-  let f : ptr<function, f32> = &(v.y);
-  *(f) = 5.0;
-}