diff --git a/src/back/spv/writer.rs b/src/back/spv/writer.rs
index d20fe48bb1..c370b0aa69 100644
--- a/src/back/spv/writer.rs
+++ b/src/back/spv/writer.rs
@@ -2960,6 +2960,11 @@ impl Writer {
         for capability in self.capabilities.iter() {
             Instruction::capability(*capability).to_words(&mut self.logical_layout.capabilities);
         }
+        if ir_module.entry_points.is_empty() {
+            // SPIR-V doesn't like modules without entry points
+            Instruction::capability(spirv::Capability::Linkage)
+                .to_words(&mut self.logical_layout.capabilities);
+        }
 
         let addressing_model = spirv::AddressingModel::Logical;
         let memory_model = spirv::MemoryModel::GLSL450;
diff --git a/src/front/wgsl/lexer.rs b/src/front/wgsl/lexer.rs
index a0d1d64963..a6a5eead39 100644
--- a/src/front/wgsl/lexer.rs
+++ b/src/front/wgsl/lexer.rs
@@ -129,7 +129,7 @@ fn consume_token(mut input: &str, generic: bool) -> (Token<'_>, &str) {
             }
         }
         '+' | '*' | '/' | '%' | '^' => (Token::Operation(cur), chars.as_str()),
-        '!' => {
+        '!' | '~' => {
             input = chars.as_str();
             if chars.next() == Some('=') {
                 (Token::LogicalOperation(cur), chars.as_str())
diff --git a/src/front/wgsl/mod.rs b/src/front/wgsl/mod.rs
index 9d30fe1af6..372c78fe02 100644
--- a/src/front/wgsl/mod.rs
+++ b/src/front/wgsl/mod.rs
@@ -1328,7 +1328,7 @@ impl Parser {
                 };
                 (true, ctx.expressions.append(expr))
             }
-            Token::Operation('!') => {
+            Token::Operation('!') | Token::Operation('~') => {
                 let expr = crate::Expression::Unary {
                     op: crate::UnaryOperator::Not,
                     expr: self.parse_singular_expression(lexer, ctx.reborrow())?,
diff --git a/src/valid/expression.rs b/src/valid/expression.rs
index 60bb169011..46d2099c6e 100644
--- a/src/valid/expression.rs
+++ b/src/valid/expression.rs
@@ -549,6 +549,7 @@ impl super::Validator {
                 match (op, inner.scalar_kind()) {
                     (_, Some(Sk::Sint))
                     | (_, Some(Sk::Bool))
+                    //TODO: restrict Negate for bools?
                     | (Uo::Negate, Some(Sk::Float))
                     | (Uo::Not, Some(Sk::Uint)) => {}
                     other => {
diff --git a/tests/in/operators.wgsl b/tests/in/operators.wgsl
index 7808e6e7c5..3b7d1f33f9 100644
--- a/tests/in/operators.wgsl
+++ b/tests/in/operators.wgsl
@@ -1,6 +1,10 @@
-[[stage(vertex)]]
-fn splat() -> [[builtin(position)]] vec4<f32> {
+fn splat() -> vec4<f32> {
     let a = (1.0 + vec2<f32>(2.0) - 3.0) / 4.0;
     let b = vec4<i32>(5) % 2;
     return a.xyxy + vec4<f32>(b);
 }
+
+fn unary() -> i32 {
+    let a = 1;
+    if (!true) { return a; } else { return ~a; };
+}
diff --git a/tests/out/operators.msl b/tests/out/operators.msl
index 4d3c3d9f30..9deefef7a0 100644
--- a/tests/out/operators.msl
+++ b/tests/out/operators.msl
@@ -2,10 +2,16 @@
 #include <simd/simd.h>
 
 
-struct splatOutput {
-    metal::float4 member [[position]];
-};
-vertex splatOutput splat(
+metal::float4 splat(
 ) {
-    return splatOutput { (((float2(1.0) + float2(2.0)) - float2(3.0)) / float2(4.0)).xyxy + static_cast<float4>(int4(5) % int4(2)) };
+    return (((float2(1.0) + float2(2.0)) - float2(3.0)) / float2(4.0)).xyxy + static_cast<float4>(int4(5) % int4(2));
+}
+
+int unary(
+) {
+    if (!true) {
+        return 1;
+    } else {
+        return !1;
+    }
 }
diff --git a/tests/out/operators.spvasm b/tests/out/operators.spvasm
index cc7c225112..a6380e59c1 100644
--- a/tests/out/operators.spvasm
+++ b/tests/out/operators.spvasm
@@ -1,12 +1,11 @@
 ; SPIR-V
 ; Version: 1.0
 ; Generator: rspirv
-; Bound: 33
+; Bound: 44
 OpCapability Shader
+OpCapability Linkage
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
-OpEntryPoint Vertex %15 "splat" %13
-OpDecorate %13 BuiltIn Position
 %2 = OpTypeVoid
 %4 = OpTypeFloat 32
 %3 = OpConstant  %4  1.0
@@ -16,29 +15,46 @@ OpDecorate %13 BuiltIn Position
 %9 = OpTypeInt 32 1
 %8 = OpConstant  %9  5
 %10 = OpConstant  %9  2
-%11 = OpTypeVector %4 4
-%14 = OpTypePointer Output %11
-%13 = OpVariable  %14  Output
-%16 = OpTypeFunction %2
-%18 = OpTypeVector %4 2
-%26 = OpTypeVector %9 4
-%15 = OpFunction  %2  None %16
-%12 = OpLabel
-OpBranch %17
-%17 = OpLabel
-%19 = OpCompositeConstruct  %18  %5 %5
-%20 = OpCompositeConstruct  %18  %3 %3
-%21 = OpFAdd  %18  %20 %19
-%22 = OpCompositeConstruct  %18  %6 %6
-%23 = OpFSub  %18  %21 %22
-%24 = OpCompositeConstruct  %18  %7 %7
-%25 = OpFDiv  %18  %23 %24
-%27 = OpCompositeConstruct  %26  %8 %8 %8 %8
-%28 = OpCompositeConstruct  %26  %10 %10 %10 %10
-%29 = OpSMod  %26  %27 %28
-%30 = OpVectorShuffle  %11  %25 %25 0 1 0 1
-%31 = OpConvertSToF  %11  %29
-%32 = OpFAdd  %11  %30 %31
-OpStore %13 %32
-OpReturn
+%11 = OpConstant  %9  1
+%13 = OpTypeBool
+%12 = OpConstantTrue  %13
+%14 = OpTypeVector %4 4
+%17 = OpTypeFunction %14
+%19 = OpTypeVector %4 2
+%27 = OpTypeVector %9 4
+%36 = OpTypeFunction %9
+%43 = OpConstantNull  %9
+%16 = OpFunction  %14  None %17
+%15 = OpLabel
+OpBranch %18
+%18 = OpLabel
+%20 = OpCompositeConstruct  %19  %5 %5
+%21 = OpCompositeConstruct  %19  %3 %3
+%22 = OpFAdd  %19  %21 %20
+%23 = OpCompositeConstruct  %19  %6 %6
+%24 = OpFSub  %19  %22 %23
+%25 = OpCompositeConstruct  %19  %7 %7
+%26 = OpFDiv  %19  %24 %25
+%28 = OpCompositeConstruct  %27  %8 %8 %8 %8
+%29 = OpCompositeConstruct  %27  %10 %10 %10 %10
+%30 = OpSMod  %27  %28 %29
+%31 = OpVectorShuffle  %14  %26 %26 0 1 0 1
+%32 = OpConvertSToF  %14  %30
+%33 = OpFAdd  %14  %31 %32
+OpReturnValue %33
+OpFunctionEnd
+%35 = OpFunction  %9  None %36
+%34 = OpLabel
+OpBranch %37
+%37 = OpLabel
+%38 = OpLogicalNot  %13  %12
+OpSelectionMerge %39 None
+OpBranchConditional %38 %40 %41
+%40 = OpLabel
+OpReturnValue %11
+%41 = OpLabel
+%42 = OpNot  %9  %11
+OpReturnValue %42
+%39 = OpLabel
+OpReturnValue %43
 OpFunctionEnd
\ No newline at end of file