[libclc] Workaround AMDGPU GlobalOffset assert

After removing the always_inline, we are getting assert in getLoads in
GlobalOffset because of phi.

Add back the always inline temporarily to workaround the assert.

Author: jsji

libclc/libspirv/lib/amdgcn/workitem/get_group_id.cl

@@ -8,7 +8,7 @@
 
 #include <libspirv/spirv.h>
 
-_CLC_DEF _CLC_OVERLOAD size_t __spirv_BuiltInWorkgroupId(int dim) {
+__attribute__((always_inline)) _CLC_DEF _CLC_OVERLOAD size_t __spirv_BuiltInWorkgroupId(int dim) {
   switch (dim) {
   case 0:
     return __builtin_amdgcn_workgroup_id_x();

libclc/libspirv/lib/generic/workitem/get_global_id.cl

@@ -8,7 +8,7 @@
 
 #include <libspirv/spirv.h>
 
-_CLC_DEF _CLC_OVERLOAD size_t __spirv_BuiltInGlobalInvocationId(int dim) {
+__attribute__((always_inline)) _CLC_DEF _CLC_OVERLOAD size_t __spirv_BuiltInGlobalInvocationId(int dim) {
   return __spirv_BuiltInWorkgroupId(dim) * __spirv_BuiltInWorkgroupSize(dim) +
          __spirv_BuiltInLocalInvocationId(dim) +
          __spirv_BuiltInGlobalOffset(dim);