Index: include/llvm/IR/IntrinsicsAMDGPU.td
===================================================================
--- include/llvm/IR/IntrinsicsAMDGPU.td
+++ include/llvm/IR/IntrinsicsAMDGPU.td
@@ -33,6 +33,8 @@
                                        "__builtin_r600_read_tgid">;
 defm int_r600_read_tidig : AMDGPUReadPreloadRegisterIntrinsic_xyz <
                                        "__builtin_r600_read_tidig">;
+defm int_r600_read_global_offset : AMDGPUReadPreloadRegisterIntrinsic_xyz <
+                                       "__builtin_r600_read_global_offset">;
 
 def int_r600_rat_store_typed :
   // 1st parameter: Data
@@ -64,6 +66,8 @@
   "__builtin_amdgcn_local_size">;
 defm int_amdgcn_global_size : AMDGPUReadPreloadRegisterIntrinsic_xyz <
   "__builtin_amdgcn_global_size">;
+defm int_amdgcn_global_offset : AMDGPUReadPreloadRegisterIntrinsic_xyz <
+  "__builtin_amdgcn_global_offset">;
 defm int_amdgcn_ngroups : AMDGPUReadPreloadRegisterIntrinsic_xyz <
   "__builtin_amdgcn_ngroups">;
 defm int_amdgcn_workitem_id : AMDGPUReadPreloadRegisterIntrinsic_xyz <
Index: lib/Target/AMDGPU/AMDGPUISelLowering.h
===================================================================
--- lib/Target/AMDGPU/AMDGPUISelLowering.h
+++ lib/Target/AMDGPU/AMDGPUISelLowering.h
@@ -208,7 +208,9 @@
 
   enum ImplicitParameter {
     GRID_DIM,
-    GRID_OFFSET
+    GRID_OFFSET_X,
+    GRID_OFFSET_Y,
+    GRID_OFFSET_Z,
   };
 
   /// \brief Helper function that returns the byte offset of the given
Index: lib/Target/AMDGPU/AMDGPUISelLowering.cpp
===================================================================
--- lib/Target/AMDGPU/AMDGPUISelLowering.cpp
+++ lib/Target/AMDGPU/AMDGPUISelLowering.cpp
@@ -2655,8 +2655,12 @@
   switch (Param) {
   case GRID_DIM:
     return ArgOffset;
-  case GRID_OFFSET:
+  case GRID_OFFSET_X:
     return ArgOffset + 4;
+  case GRID_OFFSET_Y:
+    return ArgOffset + 8;
+  case GRID_OFFSET_Z:
+    return ArgOffset + 12;
   }
   llvm_unreachable("unexpected implicit parameter type");
 }
Index: lib/Target/AMDGPU/R600ISelLowering.cpp
===================================================================
--- lib/Target/AMDGPU/R600ISelLowering.cpp
+++ lib/Target/AMDGPU/R600ISelLowering.cpp
@@ -791,6 +791,18 @@
       uint32_t ByteOffset = getImplicitParameterOffset(MFI, GRID_DIM);
       return LowerImplicitParameter(DAG, VT, DL, ByteOffset / 4);
     }
+    case Intrinsic::r600_read_global_offset_x: {
+      uint32_t ByteOffset = getImplicitParameterOffset(MFI, GRID_OFFSET_X);
+      return LowerImplicitParameter(DAG, VT, DL, ByteOffset / 4);
+    }
+    case Intrinsic::r600_read_global_offset_y: {
+      uint32_t ByteOffset = getImplicitParameterOffset(MFI, GRID_OFFSET_Y);
+      return LowerImplicitParameter(DAG, VT, DL, ByteOffset / 4);
+    }
+    case Intrinsic::r600_read_global_offset_z: {
+      uint32_t ByteOffset = getImplicitParameterOffset(MFI, GRID_OFFSET_Z);
+      return LowerImplicitParameter(DAG, VT, DL, ByteOffset / 4);
+    }
 
     case Intrinsic::r600_read_tgid_x:
       return CreateLiveInRegister(DAG, &AMDGPU::R600_TReg32RegClass,
Index: lib/Target/AMDGPU/SIISelLowering.cpp
===================================================================
--- lib/Target/AMDGPU/SIISelLowering.cpp
+++ lib/Target/AMDGPU/SIISelLowering.cpp
@@ -1657,6 +1657,28 @@
     // Really only 2 bits.
     return lowerImplicitZextParam(DAG, Op, MVT::i8,
                                   getImplicitParameterOffset(MFI, GRID_DIM));
+  case Intrinsic::amdgcn_global_offset_x:
+  case Intrinsic::r600_read_global_offset_x:
+    if (Subtarget->isAmdHsaOS())
+      return emitNonHSAIntrinsicError(DAG, VT);
+
+    return LowerParameter(DAG, VT, VT, DL, DAG.getEntryNode(),
+                        getImplicitParameterOffset(MFI, GRID_OFFSET_X), false);
+  case Intrinsic::amdgcn_global_offset_y:
+  case Intrinsic::r600_read_global_offset_y:
+    if (Subtarget->isAmdHsaOS())
+      return emitNonHSAIntrinsicError(DAG, VT);
+
+    return LowerParameter(DAG, VT, VT, DL, DAG.getEntryNode(),
+                        getImplicitParameterOffset(MFI, GRID_OFFSET_Y), false);
+  case Intrinsic::amdgcn_global_offset_z:
+  case Intrinsic::r600_read_global_offset_z:
+    if (Subtarget->isAmdHsaOS())
+      return emitNonHSAIntrinsicError(DAG, VT);
+
+    return LowerParameter(DAG, VT, VT, DL, DAG.getEntryNode(),
+                        getImplicitParameterOffset(MFI, GRID_OFFSET_Z), false);
+
   case Intrinsic::amdgcn_workgroup_id_x:
   case Intrinsic::r600_read_tgid_x:
     return CreateLiveInRegister(DAG, &AMDGPU::SReg_32RegClass,