[llvm] [SPIRV] Enable DCE in instruction selection and update tests (PR #168428)

[Steven Perron via llvm-commits]

================
@@ -506,22 +509,193 @@ static bool isConstReg(MachineRegisterInfo *MRI, Register OpReg) {
   return false;
 }
 
+static bool intrinsicHasSideEffects(Intrinsic::ID ID) {
+  switch (ID) {
+  // Intrinsics that do not have side effects.
+  // This is not an exhaustive list and may need to be updated.
+  case Intrinsic::spv_all:
+  case Intrinsic::spv_alloca:
+  case Intrinsic::spv_any:
+  case Intrinsic::spv_bitcast:
+  case Intrinsic::spv_const_composite:
+  case Intrinsic::spv_cross:
+  case Intrinsic::spv_degrees:
+  case Intrinsic::spv_distance:
+  case Intrinsic::spv_extractelt:
+  case Intrinsic::spv_extractv:
+  case Intrinsic::spv_faceforward:
+  case Intrinsic::spv_fdot:
+  case Intrinsic::spv_firstbitlow:
+  case Intrinsic::spv_firstbitshigh:
+  case Intrinsic::spv_firstbituhigh:
+  case Intrinsic::spv_frac:
+  case Intrinsic::spv_gep:
+  case Intrinsic::spv_global_offset:
+  case Intrinsic::spv_global_size:
+  case Intrinsic::spv_group_id:
+  case Intrinsic::spv_insertelt:
+  case Intrinsic::spv_insertv:
+  case Intrinsic::spv_isinf:
+  case Intrinsic::spv_isnan:
+  case Intrinsic::spv_lerp:
+  case Intrinsic::spv_length:
+  case Intrinsic::spv_normalize:
+  case Intrinsic::spv_num_subgroups:
+  case Intrinsic::spv_num_workgroups:
+  case Intrinsic::spv_ptrcast:
+  case Intrinsic::spv_radians:
+  case Intrinsic::spv_reflect:
+  case Intrinsic::spv_refract:
+  case Intrinsic::spv_resource_getpointer:
+  case Intrinsic::spv_resource_handlefrombinding:
+  case Intrinsic::spv_resource_handlefromimplicitbinding:
+  case Intrinsic::spv_resource_nonuniformindex:
+  case Intrinsic::spv_rsqrt:
+  case Intrinsic::spv_saturate:
+  case Intrinsic::spv_sdot:
+  case Intrinsic::spv_sign:
+  case Intrinsic::spv_smoothstep:
+  case Intrinsic::spv_step:
+  case Intrinsic::spv_subgroup_id:
+  case Intrinsic::spv_subgroup_local_invocation_id:
+  case Intrinsic::spv_subgroup_max_size:
+  case Intrinsic::spv_subgroup_size:
+  case Intrinsic::spv_thread_id:
+  case Intrinsic::spv_thread_id_in_group:
+  case Intrinsic::spv_udot:
+  case Intrinsic::spv_undef:
+  case Intrinsic::spv_value_md:
+  case Intrinsic::spv_workgroup_size:
+    return false;
+  default:
+    return true;
----------------
s-perron wrote:

There is already a generic way for backend to flag their intrinsic as not having side effect. When that happens, the G_INTRINSIC opcode is used instead of G_INTRINSIC_W_SIDE_EFFECTS opcode. When that happens the shared code in instruction selection will see that the opcode is trivially dead, and remove it.

However if the shared code deletes the instruction, it will not update the global registry. So we mark everything (intrinsics and opcodes) as having side effects.

If we were to try to have some type of target hook, I would not want to introduce a second way to mark an intrinsic or opcode as having no side-effect. I would want to add a hook that does backend specific cleanup when deleting an instruction.

Two reasons I don't want to add that hook:

1. In general, having a side table with information like that a non-standard design. I don't want to encourage more of it.
2. Second, we have had discussions about try to replace the global registry with an analysis. It was in the works a couple month ago. @michalpaszkowski might have more information on its current state. If that happens, then we no longer need the hook, then we have created something that is now dead. I'd prefer to keep code that will potentially become obsolete confined to the SPIRV backend.

We could add in a backend specific bit to say "this has no side effects other than the GR". I can add a TODO for that. 

https://github.com/llvm/llvm-project/pull/168428