[AMDGPU] Add __builtin_amdgcn_workgroup_size_x/y/z

The main purpose of introducing these builtins is to add a range
metadata [1, 1025) on the work group size loaded from dispatch
ptr, which cannot be done by source code.

Differential Revision: https://reviews.llvm.org/D76772