diff --git a/clang/include/clang/Basic/BuiltinsSPIRVCL.td b/clang/include/clang/Basic/BuiltinsSPIRVCL.td
index 1103a0d088e8b..10320fab34a6c 100644
--- a/clang/include/clang/Basic/BuiltinsSPIRVCL.td
+++ b/clang/include/clang/Basic/BuiltinsSPIRVCL.td
@@ -10,3 +10,6 @@ include "clang/Basic/BuiltinsSPIRVBase.td"
 
 def generic_cast_to_ptr_explicit
     : SPIRVBuiltin<"void*(void*, int)", [NoThrow, Const, CustomTypeChecking]>;
+def global_size : SPIRVBuiltin<"size_t(int)", [NoThrow, Const]>;
+def global_offset : SPIRVBuiltin<"size_t(int)", [NoThrow, Const]>;
+def subgroup_max_size : SPIRVBuiltin<"uint32_t()", [NoThrow, Const]>;
diff --git a/clang/include/clang/Basic/BuiltinsSPIRVCommon.td b/clang/include/clang/Basic/BuiltinsSPIRVCommon.td
index 17bcd0b9cb783..d2ef6f99a0502 100644
--- a/clang/include/clang/Basic/BuiltinsSPIRVCommon.td
+++ b/clang/include/clang/Basic/BuiltinsSPIRVCommon.td
@@ -8,6 +8,16 @@
 
 include "clang/Basic/BuiltinsSPIRVBase.td"
 
+def num_workgroups : SPIRVBuiltin<"size_t(int)", [NoThrow, Const]>;
+def workgroup_size : SPIRVBuiltin<"size_t(int)", [NoThrow, Const]>;
+def workgroup_id : SPIRVBuiltin<"size_t(int)", [NoThrow, Const]>;
+def local_invocation_id : SPIRVBuiltin<"size_t(int)", [NoThrow, Const]>;
+def global_invocation_id : SPIRVBuiltin<"size_t(int)", [NoThrow, Const]>;
+def subgroup_size : SPIRVBuiltin<"uint32_t()", [NoThrow, Const]>;
+def num_subgroups : SPIRVBuiltin<"uint32_t()", [NoThrow, Const]>;
+def subgroup_id : SPIRVBuiltin<"uint32_t()", [NoThrow, Const]>;
+def subgroup_local_invocation_id : SPIRVBuiltin<"uint32_t()", [NoThrow, Const]>;
+
 def distance : SPIRVBuiltin<"void(...)", [NoThrow, Const]>;
 def length : SPIRVBuiltin<"void(...)", [NoThrow, Const]>;
 def smoothstep : SPIRVBuiltin<"void(...)", [NoThrow, Const, CustomTypeChecking]>;
diff --git a/clang/lib/CodeGen/CGHLSLRuntime.cpp b/clang/lib/CodeGen/CGHLSLRuntime.cpp
index cfe9dc1192d9d..2ea3ba2d8199c 100644
--- a/clang/lib/CodeGen/CGHLSLRuntime.cpp
+++ b/clang/lib/CodeGen/CGHLSLRuntime.cpp
@@ -393,17 +393,27 @@ llvm::Value *CGHLSLRuntime::emitInputSemantic(IRBuilder<> &B,
     return B.CreateCall(FunctionCallee(GroupIndex));
   }
   if (D.hasAttr<HLSLSV_DispatchThreadIDAttr>()) {
+    llvm::Intrinsic::ID IntrinID = getThreadIdIntrinsic();
     llvm::Function *ThreadIDIntrinsic =
-        CGM.getIntrinsic(getThreadIdIntrinsic());
+        llvm::Intrinsic::isOverloaded(IntrinID)
+            ? CGM.getIntrinsic(IntrinID, {CGM.Int32Ty})
+            : CGM.getIntrinsic(IntrinID);
     return buildVectorInput(B, ThreadIDIntrinsic, Ty);
   }
   if (D.hasAttr<HLSLSV_GroupThreadIDAttr>()) {
+    llvm::Intrinsic::ID IntrinID = getGroupThreadIdIntrinsic();
     llvm::Function *GroupThreadIDIntrinsic =
-        CGM.getIntrinsic(getGroupThreadIdIntrinsic());
+        llvm::Intrinsic::isOverloaded(IntrinID)
+            ? CGM.getIntrinsic(IntrinID, {CGM.Int32Ty})
+            : CGM.getIntrinsic(IntrinID);
     return buildVectorInput(B, GroupThreadIDIntrinsic, Ty);
   }
   if (D.hasAttr<HLSLSV_GroupIDAttr>()) {
-    llvm::Function *GroupIDIntrinsic = CGM.getIntrinsic(getGroupIdIntrinsic());
+    llvm::Intrinsic::ID IntrinID = getGroupIdIntrinsic();
+    llvm::Function *GroupIDIntrinsic =
+        llvm::Intrinsic::isOverloaded(IntrinID)
+            ? CGM.getIntrinsic(IntrinID, {CGM.Int32Ty})
+            : CGM.getIntrinsic(IntrinID);
     return buildVectorInput(B, GroupIDIntrinsic, Ty);
   }
   assert(false && "Unhandled parameter attribute");
diff --git a/clang/lib/CodeGen/TargetBuiltins/SPIR.cpp b/clang/lib/CodeGen/TargetBuiltins/SPIR.cpp
index 0687485cd3f80..16243951c7bec 100644
--- a/clang/lib/CodeGen/TargetBuiltins/SPIR.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/SPIR.cpp
@@ -97,6 +97,48 @@ Value *CodeGenFunction::EmitSPIRVBuiltinExpr(unsigned BuiltinID,
     Call->addRetAttr(llvm::Attribute::AttrKind::NoUndef);
     return Call;
   }
+  case SPIRV::BI__builtin_spirv_num_workgroups:
+    return Builder.CreateIntrinsic(
+        /*ReturnType=*/getTypes().ConvertType(E->getType()),
+        Intrinsic::spv_num_workgroups,
+        ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr,
+        "spv.num.workgroups");
+  case SPIRV::BI__builtin_spirv_workgroup_size:
+    return Builder.CreateIntrinsic(
+        /*ReturnType=*/getTypes().ConvertType(E->getType()),
+        Intrinsic::spv_workgroup_size,
+        ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr,
+        "spv.workgroup.size");
+  case SPIRV::BI__builtin_spirv_workgroup_id:
+    return Builder.CreateIntrinsic(
+        /*ReturnType=*/getTypes().ConvertType(E->getType()),
+        Intrinsic::spv_group_id,
+        ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr,
+        "spv.group.id");
+  case SPIRV::BI__builtin_spirv_local_invocation_id:
+    return Builder.CreateIntrinsic(
+        /*ReturnType=*/getTypes().ConvertType(E->getType()),
+        Intrinsic::spv_thread_id_in_group,
+        ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr,
+        "spv.thread.id.in.group");
+  case SPIRV::BI__builtin_spirv_global_invocation_id:
+    return Builder.CreateIntrinsic(
+        /*ReturnType=*/getTypes().ConvertType(E->getType()),
+        Intrinsic::spv_thread_id,
+        ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr,
+        "spv.thread.id");
+  case SPIRV::BI__builtin_spirv_global_size:
+    return Builder.CreateIntrinsic(
+        /*ReturnType=*/getTypes().ConvertType(E->getType()),
+        Intrinsic::spv_global_size,
+        ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr,
+        "spv.num.workgroups");
+  case SPIRV::BI__builtin_spirv_global_offset:
+    return Builder.CreateIntrinsic(
+        /*ReturnType=*/getTypes().ConvertType(E->getType()),
+        Intrinsic::spv_global_offset,
+        ArrayRef<Value *>{EmitScalarExpr(E->getArg(0))}, nullptr,
+        "spv.global.offset");
   }
   return nullptr;
 }
diff --git a/clang/lib/Headers/__clang_spirv_builtins.h b/clang/lib/Headers/__clang_spirv_builtins.h
index e344ed52571a7..9915cdfcae7cd 100644
--- a/clang/lib/Headers/__clang_spirv_builtins.h
+++ b/clang/lib/Headers/__clang_spirv_builtins.h
@@ -16,6 +16,12 @@
 #define __SPIRV_NOEXCEPT
 #endif
 
+#pragma push_macro("__size_t")
+#pragma push_macro("__uint32_t")
+#pragma push_macro("__uint64_t")
+#define __size_t __SIZE_TYPE__
+#define __uint32_t __UINT32_TYPE__
+
 #define __SPIRV_overloadable __attribute__((overloadable))
 #define __SPIRV_convergent __attribute__((convergent))
 #define __SPIRV_inline __attribute__((always_inline))
@@ -36,13 +42,41 @@
 // to establish if we can use the builtin alias. We disable builtin altogether
 // if we do not intent to use the backend. So instead of use target macros, rely
 // on a __has_builtin test.
-#if (__has_builtin(__builtin_spirv_generic_cast_to_ptr_explicit))
+#if (__has_builtin(__builtin_spirv_num_workgroups))
 #define __SPIRV_BUILTIN_ALIAS(builtin)                                         \
   __attribute__((clang_builtin_alias(builtin)))
 #else
 #define __SPIRV_BUILTIN_ALIAS(builtin)
 #endif
 
+// Builtin IDs and sizes
+
+extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_num_workgroups) __size_t
+    __spirv_NumWorkgroups(int);
+extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_workgroup_size) __size_t
+    __spirv_WorkgroupSize(int);
+extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_workgroup_id) __size_t
+    __spirv_WorkgroupId(int);
+extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_local_invocation_id) __size_t
+    __spirv_LocalInvocationId(int);
+extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_global_invocation_id) __size_t
+    __spirv_GlobalInvocationId(int);
+
+extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_global_size) __size_t
+    __spirv_GlobalSize(int);
+extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_global_offset) __size_t
+    __spirv_GlobalOffset(int);
+extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_size) __uint32_t
+    __spirv_SubgroupSize();
+extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_max_size) __uint32_t
+    __spirv_SubgroupMaxSize();
+extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_num_subgroups) __uint32_t
+    __spirv_NumSubgroups();
+extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_id) __uint32_t
+    __spirv_SubgroupId();
+extern __SPIRV_BUILTIN_ALIAS(__builtin_spirv_subgroup_local_invocation_id)
+    __uint32_t __spirv_SubgroupLocalInvocationId();
+
 // OpGenericCastToPtrExplicit
 
 extern __SPIRV_overloadable
@@ -164,6 +198,10 @@ __spirv_GenericCastToPtr_ToPrivate(__generic const volatile void *p,
   return (__private const volatile void *)p;
 }
 
+#pragma pop_macro("__size_t")
+#pragma pop_macro("__uint32_t")
+#pragma pop_macro("__uint64_t")
+
 #undef __SPIRV_overloadable
 #undef __SPIRV_convergent
 #undef __SPIRV_inline
diff --git a/clang/test/CodeGenHLSL/semantics/DispatchThreadID.hlsl b/clang/test/CodeGenHLSL/semantics/DispatchThreadID.hlsl
index 975a7264fd3f0..7aeb877072d87 100644
--- a/clang/test/CodeGenHLSL/semantics/DispatchThreadID.hlsl
+++ b/clang/test/CodeGenHLSL/semantics/DispatchThreadID.hlsl
@@ -5,7 +5,7 @@
 
 // CHECK:       define void @foo()
 // CHECK-DXIL:  %[[#ID:]] = call i32 @llvm.[[TARGET]].thread.id(i32 0)
-// CHECK-SPIRV: %[[#ID:]] = call i32 @llvm.[[TARGET]].thread.id(i32 0)
+// CHECK-SPIRV: %[[#ID:]] = call i32 @llvm.[[TARGET]].thread.id.i32(i32 0)
 // CHECK-DXIL:       call void @{{.*}}foo{{.*}}(i32 %[[#ID]])
 // CHECK-SPIRV:      call spir_func void @{{.*}}foo{{.*}}(i32 %[[#ID]])
 [shader("compute")]
@@ -13,9 +13,11 @@
 void foo(uint Idx : SV_DispatchThreadID) {}
 
 // CHECK:       define void @bar()
-// CHECK:       %[[#ID_X:]] = call i32 @llvm.[[TARGET]].thread.id(i32 0)
+// CHECK-DXIL:  %[[#ID_X:]] = call i32 @llvm.[[TARGET]].thread.id(i32 0)
+// CHECK-SPIRV: %[[#ID_X:]] = call i32 @llvm.[[TARGET]].thread.id.i32(i32 0)
 // CHECK:       %[[#ID_X_:]] = insertelement <2 x i32> poison, i32 %[[#ID_X]], i64 0
-// CHECK:       %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].thread.id(i32 1)
+// CHECK-DXIL:  %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].thread.id(i32 1)
+// CHECK-SPIRV: %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].thread.id.i32(i32 1)
 // CHECK:       %[[#ID_XY:]] = insertelement <2 x i32> %[[#ID_X_]], i32 %[[#ID_Y]], i64 1
 // CHECK-DXIL:  call void @{{.*}}bar{{.*}}(<2 x i32> %[[#ID_XY]])
 // CHECK-SPIRV:  call spir_func void @{{.*}}bar{{.*}}(<2 x i32> %[[#ID_XY]])
diff --git a/clang/test/CodeGenHLSL/semantics/SV_GroupID.hlsl b/clang/test/CodeGenHLSL/semantics/SV_GroupID.hlsl
index 3aa054afc9045..62985f9d1e2a7 100644
--- a/clang/test/CodeGenHLSL/semantics/SV_GroupID.hlsl
+++ b/clang/test/CodeGenHLSL/semantics/SV_GroupID.hlsl
@@ -4,7 +4,8 @@
 // Make sure SV_GroupID translated into dx.group.id for directx target and spv.group.id for spirv target.
 
 // CHECK:       define void @foo()
-// CHECK:       %[[#ID:]] = call i32 @llvm.[[TARGET]].group.id(i32 0)
+// CHECK-DXIL:       %[[#ID:]] = call i32 @llvm.[[TARGET]].group.id(i32 0)
+// CHECK-SPIRV:      %[[#ID:]] = call i32 @llvm.[[TARGET]].group.id.i32(i32 0)
 // CHECK-DXIL:       call void @{{.*}}foo{{.*}}(i32 %[[#ID]])
 // CHECK-SPIRV:      call spir_func void @{{.*}}foo{{.*}}(i32 %[[#ID]])
 [shader("compute")]
@@ -12,9 +13,11 @@
 void foo(uint Idx : SV_GroupID) {}
 
 // CHECK:       define void @bar()
-// CHECK:       %[[#ID_X:]] = call i32 @llvm.[[TARGET]].group.id(i32 0)
+// CHECK-DXIL:  %[[#ID_X:]] = call i32 @llvm.[[TARGET]].group.id(i32 0)
+// CHECK-SPIRV: %[[#ID_X:]] = call i32 @llvm.[[TARGET]].group.id.i32(i32 0)
 // CHECK:       %[[#ID_X_:]] = insertelement <2 x i32> poison, i32 %[[#ID_X]], i64 0
-// CHECK:       %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].group.id(i32 1)
+// CHECK-DXIL:  %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].group.id(i32 1)
+// CHECK-SPIRV: %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].group.id.i32(i32 1)
 // CHECK:       %[[#ID_XY:]] = insertelement <2 x i32> %[[#ID_X_]], i32 %[[#ID_Y]], i64 1
 // CHECK-DXIL:  call void @{{.*}}bar{{.*}}(<2 x i32> %[[#ID_XY]])
 // CHECK-SPIRV:  call spir_func void @{{.*}}bar{{.*}}(<2 x i32> %[[#ID_XY]])
@@ -23,11 +26,14 @@ void foo(uint Idx : SV_GroupID) {}
 void bar(uint2 Idx : SV_GroupID) {}
 
 // CHECK:  define void @test()
-// CHECK:  %[[#ID_X:]] = call i32 @llvm.[[TARGET]].group.id(i32 0)
+// CHECK-DXIL:  %[[#ID_X:]] = call i32 @llvm.[[TARGET]].group.id(i32 0)
+// CHECK-SPIRV: %[[#ID_X:]] = call i32 @llvm.[[TARGET]].group.id.i32(i32 0)
 // CHECK:  %[[#ID_X_:]] = insertelement <3 x i32> poison, i32 %[[#ID_X]], i64 0
-// CHECK:  %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].group.id(i32 1)
+// CHECK-DXIL:  %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].group.id(i32 1)
+// CHECK-SPIRV: %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].group.id.i32(i32 1)
 // CHECK:  %[[#ID_XY:]] = insertelement <3 x i32> %[[#ID_X_]], i32 %[[#ID_Y]], i64 1
-// CHECK:  %[[#ID_Z:]] = call i32 @llvm.[[TARGET]].group.id(i32 2)
+// CHECK-DXIL:  %[[#ID_Z:]] = call i32 @llvm.[[TARGET]].group.id(i32 2)
+// CHECK-SPIRV: %[[#ID_Z:]] = call i32 @llvm.[[TARGET]].group.id.i32(i32 2)
 // CHECK:  %[[#ID_XYZ:]] = insertelement <3 x i32> %[[#ID_XY]], i32 %[[#ID_Z]], i64 2
 // CHECK-DXIL:   call void @{{.*}}test{{.*}}(<3 x i32> %[[#ID_XYZ]])
 // CHECK-SPIRV:  call spir_func void @{{.*}}test{{.*}}(<3 x i32> %[[#ID_XYZ]])
diff --git a/clang/test/CodeGenHLSL/semantics/SV_GroupThreadID.hlsl b/clang/test/CodeGenHLSL/semantics/SV_GroupThreadID.hlsl
index 3d347b973f39c..2675c973b531a 100644
--- a/clang/test/CodeGenHLSL/semantics/SV_GroupThreadID.hlsl
+++ b/clang/test/CodeGenHLSL/semantics/SV_GroupThreadID.hlsl
@@ -4,7 +4,8 @@
 // Make sure SV_GroupThreadID translated into dx.thread.id.in.group for directx target and spv.thread.id.in.group for spirv target.
 
 // CHECK:       define void @foo()
-// CHECK:       %[[#ID:]] = call i32 @llvm.[[TARGET]].thread.id.in.group(i32 0)
+// CHECK-DXIL:       %[[#ID:]] = call i32 @llvm.[[TARGET]].thread.id.in.group(i32 0)
+// CHECK-SPIRV:      %[[#ID:]] = call i32 @llvm.[[TARGET]].thread.id.in.group.i32(i32 0)
 // CHECK-DXIL:       call void @{{.*}}foo{{.*}}(i32 %[[#ID]])
 // CHECK-SPIRV:      call spir_func void @{{.*}}foo{{.*}}(i32 %[[#ID]])
 [shader("compute")]
@@ -12,9 +13,11 @@
 void foo(uint Idx : SV_GroupThreadID) {}
 
 // CHECK:       define void @bar()
-// CHECK:       %[[#ID_X:]] = call i32 @llvm.[[TARGET]].thread.id.in.group(i32 0)
+// CHECK-DXIL:  %[[#ID_X:]] = call i32 @llvm.[[TARGET]].thread.id.in.group(i32 0)
+// CHECK-SPIRV: %[[#ID_X:]] = call i32 @llvm.[[TARGET]].thread.id.in.group.i32(i32 0)
 // CHECK:       %[[#ID_X_:]] = insertelement <2 x i32> poison, i32 %[[#ID_X]], i64 0
-// CHECK:       %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].thread.id.in.group(i32 1)
+// CHECK-DXIL:  %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].thread.id.in.group(i32 1)
+// CHECK-SPIRV: %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].thread.id.in.group.i32(i32 1)
 // CHECK:       %[[#ID_XY:]] = insertelement <2 x i32> %[[#ID_X_]], i32 %[[#ID_Y]], i64 1
 // CHECK-DXIL:  call void @{{.*}}bar{{.*}}(<2 x i32> %[[#ID_XY]])
 // CHECK-SPIRV:  call spir_func void @{{.*}}bar{{.*}}(<2 x i32> %[[#ID_XY]])
@@ -23,11 +26,14 @@ void foo(uint Idx : SV_GroupThreadID) {}
 void bar(uint2 Idx : SV_GroupThreadID) {}
 
 // CHECK:  define void @test()
-// CHECK:  %[[#ID_X:]] = call i32 @llvm.[[TARGET]].thread.id.in.group(i32 0)
+// CHECK-DXIL:  %[[#ID_X:]] = call i32 @llvm.[[TARGET]].thread.id.in.group(i32 0)
+// CHECK-SPIRV: %[[#ID_X:]] = call i32 @llvm.[[TARGET]].thread.id.in.group.i32(i32 0)
 // CHECK:  %[[#ID_X_:]] = insertelement <3 x i32> poison, i32 %[[#ID_X]], i64 0
-// CHECK:  %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].thread.id.in.group(i32 1)
+// CHECK-DXIL:  %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].thread.id.in.group(i32 1)
+// CHECK-SPIRV: %[[#ID_Y:]] = call i32 @llvm.[[TARGET]].thread.id.in.group.i32(i32 1)
 // CHECK:  %[[#ID_XY:]] = insertelement <3 x i32> %[[#ID_X_]], i32 %[[#ID_Y]], i64 1
-// CHECK:  %[[#ID_Z:]] = call i32 @llvm.[[TARGET]].thread.id.in.group(i32 2)
+// CHECK-DXIL:  %[[#ID_Z:]] = call i32 @llvm.[[TARGET]].thread.id.in.group(i32 2)
+// CHECK-SPIRV: %[[#ID_Z:]] = call i32 @llvm.[[TARGET]].thread.id.in.group.i32(i32 2)
 // CHECK:  %[[#ID_XYZ:]] = insertelement <3 x i32> %[[#ID_XY]], i32 %[[#ID_Z]], i64 2
 // CHECK-DXIL:   call void @{{.*}}test{{.*}}(<3 x i32> %[[#ID_XYZ]])
 // CHECK-SPIRV:  call spir_func void @{{.*}}test{{.*}}(<3 x i32> %[[#ID_XYZ]])
diff --git a/clang/test/CodeGenSPIRV/Builtins/ids_and_ranges.c b/clang/test/CodeGenSPIRV/Builtins/ids_and_ranges.c
new file mode 100644
index 0000000000000..f71af779ec358
--- /dev/null
+++ b/clang/test/CodeGenSPIRV/Builtins/ids_and_ranges.c
@@ -0,0 +1,106 @@
+// RUN: %clang_cc1 -O1 -triple spirv64 -fsycl-is-device %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK64
+// RUN: %clang_cc1 -O1 -triple spirv64 -cl-std=CL3.0 -x cl %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK64
+// RUN: %clang_cc1 -O1 -triple spirv32 -cl-std=CL3.0 -x cl %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK32
+
+// CHECK: @test_num_workgroups(
+// CHECK-NEXT:  [[ENTRY:.*:]]
+// CHECK64-NEXT:    tail call i64 @llvm.spv.num.workgroups.i64(i32 0)
+// CHECK32-NEXT:    tail call i32 @llvm.spv.num.workgroups.i32(i32 0)
+//
+unsigned int test_num_workgroups() {
+    return __builtin_spirv_num_workgroups(0);
+}
+
+// CHECK: @test_workgroup_size(
+// CHECK-NEXT:  [[ENTRY:.*:]]
+// CHECK64-NEXT:    tail call i64 @llvm.spv.workgroup.size.i64(i32 0)
+// CHECK32-NEXT:    tail call i32 @llvm.spv.workgroup.size.i32(i32 0)
+//
+unsigned int test_workgroup_size() {
+    return __builtin_spirv_workgroup_size(0);
+}
+
+// CHECK: @test_workgroup_id(
+// CHECK-NEXT:  [[ENTRY:.*:]]
+// CHECK64-NEXT:    tail call i64 @llvm.spv.group.id.i64(i32 0)
+// CHECK32-NEXT:    tail call i32 @llvm.spv.group.id.i32(i32 0)
+//
+unsigned int test_workgroup_id() {
+    return __builtin_spirv_workgroup_id(0);
+}
+
+// CHECK: @test_local_invocation_id(
+// CHECK-NEXT:  [[ENTRY:.*:]]
+// CHECK64-NEXT:    tail call i64 @llvm.spv.thread.id.in.group.i64(i32 0)
+// CHECK32-NEXT:    tail call i32 @llvm.spv.thread.id.in.group.i32(i32 0)
+//
+unsigned int test_local_invocation_id() {
+    return __builtin_spirv_local_invocation_id(0);
+}
+
+// CHECK: @test_global_invocation_id(
+// CHECK-NEXT:  [[ENTRY:.*:]]
+// CHECK64-NEXT:    tail call i64 @llvm.spv.thread.id.i64(i32 0)
+// CHECK32-NEXT:    tail call i32 @llvm.spv.thread.id.i32(i32 0)
+//
+unsigned int test_global_invocation_id() {
+    return __builtin_spirv_global_invocation_id(0);
+}
+
+// CHECK: @test_global_size(
+// CHECK-NEXT:  [[ENTRY:.*:]]
+// CHECK64-NEXT:    tail call i64 @llvm.spv.global.size.i64(i32 0)
+// CHECK32-NEXT:    tail call i32 @llvm.spv.global.size.i32(i32 0)
+//
+unsigned int test_global_size() {
+    return __builtin_spirv_global_size(0);
+}
+
+// CHECK: @test_global_offset(
+// CHECK-NEXT:  [[ENTRY:.*:]]
+// CHECK64-NEXT:    tail call i64 @llvm.spv.global.offset.i64(i32 0)
+// CHECK32-NEXT:    tail call i32 @llvm.spv.global.offset.i32(i32 0)
+//
+unsigned int test_global_offset() {
+    return __builtin_spirv_global_offset(0);
+}
+
+// CHECK: @test_subgroup_size(
+// CHECK-NEXT:  [[ENTRY:.*:]]
+// CHECK-NEXT:    tail call i32 @llvm.spv.subgroup.size()
+//
+unsigned int test_subgroup_size() {
+    return __builtin_spirv_subgroup_size();
+}
+
+// CHECK: @test_subgroup_max_size(
+// CHECK-NEXT:  [[ENTRY:.*:]]
+// CHECK-NEXT:    tail call i32 @llvm.spv.subgroup.max.size()
+//
+unsigned int test_subgroup_max_size() {
+    return __builtin_spirv_subgroup_max_size();
+}
+
+// CHECK: @test_num_subgroups(
+// CHECK-NEXT:  [[ENTRY:.*:]]
+// CHECK-NEXT:    tail call i32 @llvm.spv.num.subgroups()
+//
+unsigned int test_num_subgroups() {
+    return __builtin_spirv_num_subgroups();
+}
+
+// CHECK: @test_subgroup_id(
+// CHECK-NEXT:  [[ENTRY:.*:]]
+// CHECK-NEXT:    tail call i32 @llvm.spv.subgroup.id()
+//
+unsigned int test_subgroup_id() {
+    return __builtin_spirv_subgroup_id();
+}
+
+// CHECK: @test_subgroup_local_invocation_id(
+// CHECK-NEXT:  [[ENTRY:.*:]]
+// CHECK-NEXT:    tail call i32 @llvm.spv.subgroup.local.invocation.id()
+//
+unsigned int test_subgroup_local_invocation_id() {
+    return __builtin_spirv_subgroup_local_invocation_id();
+}
diff --git a/clang/test/Headers/spirv_ids.cpp b/clang/test/Headers/spirv_ids.cpp
new file mode 100644
index 0000000000000..0cd74dbca53aa
--- /dev/null
+++ b/clang/test/Headers/spirv_ids.cpp
@@ -0,0 +1,110 @@
+// RUN: %clang_cc1 -Wno-unused-value -O0 -internal-isystem %S/../../lib/Headers -include __clang_spirv_builtins.h -triple spirv64 -emit-llvm %s -fsycl-is-device -o - | FileCheck %s -check-prefixes=CHECK64
+// RUN: %clang_cc1 -Wno-unused-value -O0 -internal-isystem %S/../../lib/Headers -include __clang_spirv_builtins.h -triple spirv64 -emit-llvm %s -x cl -o - | FileCheck %s -check-prefixes=CHECK64
+// RUN: %clang_cc1 -Wno-unused-value -O0 -internal-isystem %S/../../lib/Headers -include __clang_spirv_builtins.h -triple spirv32 -emit-llvm %s -fsycl-is-device -o - | FileCheck %s -check-prefixes=CHECK32
+// RUN: %clang_cc1 -Wno-unused-value -O0 -internal-isystem %S/../../lib/Headers -include __clang_spirv_builtins.h -triple spirv32 -emit-llvm %s -x cl -o - | FileCheck %s -check-prefixes=CHECK32
+// RUN: %clang_cc1 -Wno-unused-value -O0 -internal-isystem %S/../../lib/Headers -include __clang_spirv_builtins.h -triple nvptx64 -emit-llvm %s -fsycl-is-device -o - | FileCheck %s -check-prefixes=NV
+
+
+// CHECK64: call i64 @llvm.spv.num.workgroups.i64(i32 0)
+// CHECK64: call i64 @llvm.spv.num.workgroups.i64(i32 1)
+// CHECK64: call i64 @llvm.spv.num.workgroups.i64(i32 2)
+// CHECK64: call i64 @llvm.spv.workgroup.size.i64(i32 0)
+// CHECK64: call i64 @llvm.spv.workgroup.size.i64(i32 1)
+// CHECK64: call i64 @llvm.spv.workgroup.size.i64(i32 2)
+// CHECK64: call i64 @llvm.spv.group.id.i64(i32 0)
+// CHECK64: call i64 @llvm.spv.group.id.i64(i32 1)
+// CHECK64: call i64 @llvm.spv.group.id.i64(i32 2)
+// CHECK64: call i64 @llvm.spv.thread.id.in.group.i64(i32 0)
+// CHECK64: call i64 @llvm.spv.thread.id.in.group.i64(i32 1)
+// CHECK64: call i64 @llvm.spv.thread.id.in.group.i64(i32 2)
+// CHECK64: call i64 @llvm.spv.thread.id.i64(i32 0)
+// CHECK64: call i64 @llvm.spv.thread.id.i64(i32 1)
+// CHECK64: call i64 @llvm.spv.thread.id.i64(i32 2)
+// CHECK64: call i64 @llvm.spv.global.size.i64(i32 0)
+// CHECK64: call i64 @llvm.spv.global.size.i64(i32 1)
+// CHECK64: call i64 @llvm.spv.global.size.i64(i32 2)
+// CHECK64: call i64 @llvm.spv.global.offset.i64(i32 0)
+// CHECK64: call i64 @llvm.spv.global.offset.i64(i32 1)
+// CHECK64: call i64 @llvm.spv.global.offset.i64(i32 2)
+// CHECK32: call i32 @llvm.spv.num.workgroups.i32(i32 0)
+// CHECK32: call i32 @llvm.spv.num.workgroups.i32(i32 1)
+// CHECK32: call i32 @llvm.spv.num.workgroups.i32(i32 2)
+// CHECK32: call i32 @llvm.spv.workgroup.size.i32(i32 0)
+// CHECK32: call i32 @llvm.spv.workgroup.size.i32(i32 1)
+// CHECK32: call i32 @llvm.spv.workgroup.size.i32(i32 2)
+// CHECK32: call i32 @llvm.spv.group.id.i32(i32 0)
+// CHECK32: call i32 @llvm.spv.group.id.i32(i32 1)
+// CHECK32: call i32 @llvm.spv.group.id.i32(i32 2)
+// CHECK32: call i32 @llvm.spv.thread.id.in.group.i32(i32 0)
+// CHECK32: call i32 @llvm.spv.thread.id.in.group.i32(i32 1)
+// CHECK32: call i32 @llvm.spv.thread.id.in.group.i32(i32 2)
+// CHECK32: call i32 @llvm.spv.thread.id.i32(i32 0)
+// CHECK32: call i32 @llvm.spv.thread.id.i32(i32 1)
+// CHECK32: call i32 @llvm.spv.thread.id.i32(i32 2)
+// CHECK32: call i32 @llvm.spv.global.size.i32(i32 0)
+// CHECK32: call i32 @llvm.spv.global.size.i32(i32 1)
+// CHECK32: call i32 @llvm.spv.global.size.i32(i32 2)
+// CHECK32: call i32 @llvm.spv.global.offset.i32(i32 0)
+// CHECK32: call i32 @llvm.spv.global.offset.i32(i32 1)
+// CHECK32: call i32 @llvm.spv.global.offset.i32(i32 2)
+// CHECK: call i32 @llvm.spv.subgroup.size()
+// CHECK: call i32 @llvm.spv.subgroup.max.size()
+// CHECK: call i32 @llvm.spv.num.subgroups()
+// CHECK: call i32 @llvm.spv.subgroup.id()
+// CHECK: call i32 @llvm.spv.subgroup.local.invocation.id()
+  
+// NV: call noundef i64 @_Z21__spirv_NumWorkgroupsi(i32 noundef 0) #2
+// NV: call noundef i64 @_Z21__spirv_NumWorkgroupsi(i32 noundef 1) #2
+// NV: call noundef i64 @_Z21__spirv_NumWorkgroupsi(i32 noundef 2) #2
+// NV: call noundef i64 @_Z21__spirv_WorkgroupSizei(i32 noundef 0) #2
+// NV: call noundef i64 @_Z21__spirv_WorkgroupSizei(i32 noundef 1) #2
+// NV: call noundef i64 @_Z21__spirv_WorkgroupSizei(i32 noundef 2) #2
+// NV: call noundef i64 @_Z19__spirv_WorkgroupIdi(i32 noundef 0) #2
+// NV: call noundef i64 @_Z19__spirv_WorkgroupIdi(i32 noundef 1) #2
+// NV: call noundef i64 @_Z19__spirv_WorkgroupIdi(i32 noundef 2) #2
+// NV: call noundef i64 @_Z25__spirv_LocalInvocationIdi(i32 noundef 0) #2
+// NV: call noundef i64 @_Z25__spirv_LocalInvocationIdi(i32 noundef 1) #2
+// NV: call noundef i64 @_Z25__spirv_LocalInvocationIdi(i32 noundef 2) #2
+// NV: call noundef i64 @_Z26__spirv_GlobalInvocationIdi(i32 noundef 0) #2
+// NV: call noundef i64 @_Z26__spirv_GlobalInvocationIdi(i32 noundef 1) #2
+// NV: call noundef i64 @_Z26__spirv_GlobalInvocationIdi(i32 noundef 2) #2
+// NV: call noundef i64 @_Z18__spirv_GlobalSizei(i32 noundef 0) #2
+// NV: call noundef i64 @_Z18__spirv_GlobalSizei(i32 noundef 1) #2
+// NV: call noundef i64 @_Z18__spirv_GlobalSizei(i32 noundef 2) #2
+// NV: call noundef i64 @_Z20__spirv_GlobalOffseti(i32 noundef 0) #2
+// NV: call noundef i64 @_Z20__spirv_GlobalOffseti(i32 noundef 1) #2
+// NV: call noundef i64 @_Z20__spirv_GlobalOffseti(i32 noundef 2) #2
+// NV: call noundef i32 @_Z20__spirv_SubgroupSizev() #2
+// NV: call noundef i32 @_Z23__spirv_SubgroupMaxSizev() #2
+// NV: call noundef i32 @_Z20__spirv_NumSubgroupsv() #2
+// NV: call noundef i32 @_Z18__spirv_SubgroupIdv() #2
+// NV: call noundef i32 @_Z33__spirv_SubgroupLocalInvocationIdv() #2
+
+void test_id_and_range() {
+  __spirv_NumWorkgroups(0);
+  __spirv_NumWorkgroups(1);
+  __spirv_NumWorkgroups(2);
+  __spirv_WorkgroupSize(0);
+  __spirv_WorkgroupSize(1);
+  __spirv_WorkgroupSize(2);
+  __spirv_WorkgroupId(0);
+  __spirv_WorkgroupId(1);
+  __spirv_WorkgroupId(2);
+  __spirv_LocalInvocationId(0);
+  __spirv_LocalInvocationId(1);
+  __spirv_LocalInvocationId(2);
+  __spirv_GlobalInvocationId(0);
+  __spirv_GlobalInvocationId(1);
+  __spirv_GlobalInvocationId(2);
+  __spirv_GlobalSize(0);
+  __spirv_GlobalSize(1);
+  __spirv_GlobalSize(2);
+  __spirv_GlobalOffset(0);
+  __spirv_GlobalOffset(1);
+  __spirv_GlobalOffset(2);
+  unsigned int ssize = __spirv_SubgroupSize();
+  unsigned int smax = __spirv_SubgroupMaxSize();
+  unsigned int snum = __spirv_NumSubgroups();
+  unsigned int sid = __spirv_SubgroupId();
+  unsigned int sinvocid = __spirv_SubgroupLocalInvocationId();
+}
diff --git a/clang/test/SemaSPIRV/BuiltIns/ids_and_ranges.c b/clang/test/SemaSPIRV/BuiltIns/ids_and_ranges.c
new file mode 100644
index 0000000000000..0d98a552bb1b9
--- /dev/null
+++ b/clang/test/SemaSPIRV/BuiltIns/ids_and_ranges.c
@@ -0,0 +1,77 @@
+// RUN: %clang_cc1 -O1 -Wno-unused-value -triple spirv64 -fsycl-is-device -verify %s -o -
+// RUN: %clang_cc1 -O1 -Wno-unused-value -triple spirv64 -verify %s -cl-std=CL3.0 -x cl -o -
+// RUN: %clang_cc1 -O1 -Wno-unused-value -triple spirv32 -verify %s -cl-std=CL3.0 -x cl -o -
+
+void test_num_workgroups(int* p) {
+  __builtin_spirv_num_workgroups(0);
+  __builtin_spirv_num_workgroups(p); // expected-error{{incompatible pointer to integer conversion}}
+  __builtin_spirv_num_workgroups(0, 0); // expected-error{{too many arguments to function call, expected 1, have 2}}
+  __builtin_spirv_num_workgroups(); // expected-error{{too few arguments to function call, expected 1, have 0}}
+}
+
+void test_workgroup_size(int* p) {
+  __builtin_spirv_workgroup_size(0);
+  __builtin_spirv_workgroup_size(p); // expected-error{{incompatible pointer to integer conversion}}
+  __builtin_spirv_workgroup_size(0, 0); // expected-error{{too many arguments to function call, expected 1, have 2}}
+  __builtin_spirv_workgroup_size(); // expected-error{{too few arguments to function call, expected 1, have 0}}
+}
+
+void test_workgroup_id(int* p) {
+  __builtin_spirv_workgroup_id(0);
+  __builtin_spirv_workgroup_id(p); // expected-error{{incompatible pointer to integer conversion}}
+  __builtin_spirv_workgroup_id(0, 0); // expected-error{{too many arguments to function call, expected 1, have 2}}
+  __builtin_spirv_workgroup_id(); // expected-error{{too few arguments to function call, expected 1, have 0}}
+}
+
+void test_local_invocation_id(int* p) {
+  __builtin_spirv_local_invocation_id(0);
+  __builtin_spirv_local_invocation_id(p); // expected-error{{incompatible pointer to integer conversion}}
+  __builtin_spirv_local_invocation_id(0, 0); // expected-error{{too many arguments to function call, expected 1, have 2}}
+  __builtin_spirv_local_invocation_id(); // expected-error{{too few arguments to function call, expected 1, have 0}}
+}
+
+void test_global_invocation_id(int* p) {
+  __builtin_spirv_global_invocation_id(0);
+  __builtin_spirv_global_invocation_id(p); // expected-error{{incompatible pointer to integer conversion}}
+  __builtin_spirv_global_invocation_id(0, 0); // expected-error{{too many arguments to function call, expected 1, have 2}}
+  __builtin_spirv_global_invocation_id(); // expected-error{{too few arguments to function call, expected 1, have 0}}
+}
+
+void test_global_size(int* p) {
+  __builtin_spirv_global_size(0);
+  __builtin_spirv_global_size(p); // expected-error{{incompatible pointer to integer conversion}}
+  __builtin_spirv_global_size(0, 0); // expected-error{{too many arguments to function call, expected 1, have 2}}
+  __builtin_spirv_global_size(); // expected-error{{too few arguments to function call, expected 1, have 0}}
+}
+
+void test_global_offset(int* p) {
+  __builtin_spirv_global_offset(0);
+  __builtin_spirv_global_offset(p); // expected-error{{incompatible pointer to integer conversion}}
+  __builtin_spirv_global_offset(0, 0); // expected-error{{too many arguments to function call, expected 1, have 2}}
+  __builtin_spirv_global_offset(); // expected-error{{too few arguments to function call, expected 1, have 0}}
+}
+
+void test_subgroup_size() {
+  __builtin_spirv_subgroup_size();
+  __builtin_spirv_subgroup_size(0); // expected-error{{too many arguments to function call, expected 0, have 1}}
+}
+
+void test_subgroup_max_size() {
+  __builtin_spirv_subgroup_max_size();
+  __builtin_spirv_subgroup_max_size(0); // expected-error{{too many arguments to function call, expected 0, have 1}}
+}
+
+void test_num_subgroups() {
+  __builtin_spirv_num_subgroups();
+  __builtin_spirv_num_subgroups(0); // expected-error{{too many arguments to function call, expected 0, have 1}}
+}
+
+void test_subgroup_id() {
+  __builtin_spirv_subgroup_id();
+  __builtin_spirv_subgroup_id(0); // expected-error{{too many arguments to function call, expected 0, have 1}}
+}
+
+void test_subgroup_local_invocation_id() {
+  __builtin_spirv_subgroup_local_invocation_id();
+  __builtin_spirv_subgroup_local_invocation_id(0); // expected-error{{too many arguments to function call, expected 0, have 1}}
+}
diff --git a/llvm/include/llvm/IR/IntrinsicsSPIRV.td b/llvm/include/llvm/IR/IntrinsicsSPIRV.td
index 8d984d6ce58df..a60252f6e0886 100644
--- a/llvm/include/llvm/IR/IntrinsicsSPIRV.td
+++ b/llvm/include/llvm/IR/IntrinsicsSPIRV.td
@@ -59,10 +59,24 @@ let TargetPrefix = "spv" in {
                                 NoCapture<ArgIndex<1>>,
                                 ImmArg<ArgIndex<0>>]>;
 
-  // The following intrinsic(s) are mirrored from IntrinsicsDirectX.td for HLSL support.
-  def int_spv_thread_id : Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem, IntrWillReturn]>;
-  def int_spv_group_id : Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem, IntrWillReturn]>;
-  def int_spv_thread_id_in_group : Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem, IntrWillReturn]>;
+  // Ideally we should use the SPIR-V terminology for SPIR-V intrinsics.
+  def int_spv_thread_id : Intrinsic<[llvm_anyint_ty], [llvm_i32_ty], [IntrNoMem, IntrWillReturn]>;
+  def int_spv_group_id : Intrinsic<[llvm_anyint_ty], [llvm_i32_ty], [IntrNoMem, IntrWillReturn]>;
+  def int_spv_thread_id_in_group : Intrinsic<[llvm_anyint_ty], [llvm_i32_ty], [IntrNoMem, IntrWillReturn]>;
+  def int_spv_workgroup_size : Intrinsic<[llvm_anyint_ty], [llvm_i32_ty], [IntrNoMem, IntrWillReturn]>;
+  def int_spv_global_size : Intrinsic<[llvm_anyint_ty], [llvm_i32_ty], [IntrNoMem, IntrWillReturn]>;
+  def int_spv_global_offset : Intrinsic<[llvm_anyint_ty], [llvm_i32_ty], [IntrNoMem, IntrWillReturn]>;
+  def int_spv_num_workgroups : Intrinsic<[llvm_anyint_ty], [llvm_i32_ty], [IntrNoMem, IntrWillReturn]>;
+  def int_spv_subgroup_size : ClangBuiltin<"__builtin_spirv_subgroup_size">,
+                          Intrinsic<[llvm_i32_ty], [], [NoUndef<RetIndex>, IntrNoMem, IntrWillReturn]>;
+  def int_spv_num_subgroups : ClangBuiltin<"__builtin_spirv_num_subgroups">,
+                          Intrinsic<[llvm_i32_ty], [], [NoUndef<RetIndex>, IntrNoMem, IntrWillReturn]>;
+  def int_spv_subgroup_id : ClangBuiltin<"__builtin_spirv_subgroup_id">,
+                          Intrinsic<[llvm_i32_ty], [], [NoUndef<RetIndex>, IntrNoMem, IntrWillReturn]>;
+  def int_spv_subgroup_local_invocation_id : ClangBuiltin<"__builtin_spirv_subgroup_local_invocation_id">,
+                          Intrinsic<[llvm_i32_ty], [], [NoUndef<RetIndex>, IntrNoMem, IntrWillReturn]>;
+  def int_spv_subgroup_max_size : ClangBuiltin<"__builtin_spirv_subgroup_max_size">,
+                          Intrinsic<[llvm_i32_ty], [], [NoUndef<RetIndex>, IntrNoMem, IntrWillReturn]>;
   def int_spv_flattened_thread_id_in_group : Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrWillReturn]>;
   def int_spv_all : DefaultAttrsIntrinsic<[llvm_i1_ty], [llvm_any_ty], [IntrNoMem]>;
   def int_spv_any : DefaultAttrsIntrinsic<[llvm_i1_ty], [llvm_any_ty], [IntrNoMem]>;
diff --git a/llvm/lib/IR/Intrinsics.cpp b/llvm/lib/IR/Intrinsics.cpp
index e631419d5e1c2..d2632d50dff06 100644
--- a/llvm/lib/IR/Intrinsics.cpp
+++ b/llvm/lib/IR/Intrinsics.cpp
@@ -27,6 +27,7 @@
 #include "llvm/IR/IntrinsicsR600.h"
 #include "llvm/IR/IntrinsicsRISCV.h"
 #include "llvm/IR/IntrinsicsS390.h"
+#include "llvm/IR/IntrinsicsSPIRV.h"
 #include "llvm/IR/IntrinsicsVE.h"
 #include "llvm/IR/IntrinsicsX86.h"
 #include "llvm/IR/IntrinsicsXCore.h"
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
index 5258f07d2f71b..b1e14769eaf9b 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -3043,6 +3043,32 @@ bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
     // a `LocalInvocationIndex` builtin variable
     return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
                               ResType, I);
+  case Intrinsic::spv_workgroup_size:
+    return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
+                                  ResType, I);
+  case Intrinsic::spv_global_size:
+    return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
+                                  I);
+  case Intrinsic::spv_global_offset:
+    return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
+                                  ResType, I);
+  case Intrinsic::spv_num_workgroups:
+    return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
+                                  ResType, I);
+  case Intrinsic::spv_subgroup_size:
+    return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
+                              I);
+  case Intrinsic::spv_num_subgroups:
+    return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
+                              I);
+  case Intrinsic::spv_subgroup_id:
+    return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType, I);
+  case Intrinsic::spv_subgroup_local_invocation_id:
+    return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
+                              ResVReg, ResType, I);
+  case Intrinsic::spv_subgroup_max_size:
+    return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
+                              I);
   case Intrinsic::spv_fdot:
     return selectFloatDot(ResVReg, ResType, I);
   case Intrinsic::spv_udot:
@@ -3970,13 +3996,13 @@ bool SPIRVInstructionSelector::selectLog10(Register ResVReg,
 // Generate the instructions to load 3-element vector builtin input
 // IDs/Indices.
 // Like: GlobalInvocationId, LocalInvocationId, etc....
+
 bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
     SPIRV::BuiltIn::BuiltIn BuiltInValue, Register ResVReg,
     const SPIRVType *ResType, MachineInstr &I) const {
   MachineIRBuilder MIRBuilder(I);
-  const SPIRVType *U32Type = GR.getOrCreateSPIRVIntegerType(32, MIRBuilder);
   const SPIRVType *Vec3Ty =
-      GR.getOrCreateSPIRVVectorType(U32Type, 3, MIRBuilder, false);
+      GR.getOrCreateSPIRVVectorType(ResType, 3, MIRBuilder, false);
   const SPIRVType *PtrType = GR.getOrCreateSPIRVPointerType(
       Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
 
diff --git a/llvm/test/CodeGen/SPIRV/builtin_intrinsics_32.ll b/llvm/test/CodeGen/SPIRV/builtin_intrinsics_32.ll
new file mode 100644
index 0000000000000..39a755e736081
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/builtin_intrinsics_32.ll
@@ -0,0 +1,136 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1"
+target triple = "spirv32-unknown-unknown"
+
+; CHECK: OpDecorate [[NumWorkgroups:%[0-9]*]] BuiltIn NumWorkgroups
+; CHECK: OpDecorate [[WorkgroupSize:%[0-9]*]] BuiltIn WorkgroupSize
+; CHECK: OpDecorate [[WorkgroupId:%[0-9]*]] BuiltIn WorkgroupId
+; CHECK: OpDecorate [[LocalInvocationId:%[0-9]*]] BuiltIn LocalInvocationId
+; CHECK: OpDecorate [[GlobalInvocationId:%[0-9]*]] BuiltIn GlobalInvocationId
+; CHECK: OpDecorate [[GlobalSize:%[0-9]*]] BuiltIn GlobalSize
+; CHECK: OpDecorate [[GlobalOffset:%[0-9]*]] BuiltIn GlobalOffset
+; CHECK: OpDecorate [[SubgroupSize:%[0-9]*]] BuiltIn SubgroupSize
+; CHECK: OpDecorate [[SubgroupMaxSize:%[0-9]*]] BuiltIn SubgroupMaxSize
+; CHECK: OpDecorate [[NumSubgroups:%[0-9]*]] BuiltIn NumSubgroups
+; CHECK: OpDecorate [[SubgroupId:%[0-9]*]] BuiltIn SubgroupId
+; CHECK: OpDecorate [[SubgroupLocalInvocationId:%[0-9]*]] BuiltIn SubgroupLocalInvocationId
+; CHECK: [[I32:%[0-9]*]] = OpTypeInt 32 0
+; CHECK: [[I32PTR:%[0-9]*]] = OpTypePointer Input [[I32]]
+; CHECK: [[I32V3:%[0-9]*]] = OpTypeVector [[I32]] 3
+; CHECK: [[I32V3PTR:%[0-9]*]] = OpTypePointer Input [[I32V3]]
+; CHECK: [[NumWorkgroups]] = OpVariable [[I32V3PTR]] Input
+; CHECK: [[WorkgroupSize]] = OpVariable [[I32V3PTR]] Input
+; CHECK: [[WorkgroupId]] = OpVariable [[I32V3PTR]] Input
+; CHECK: [[LocalInvocationId]] = OpVariable [[I32V3PTR]] Input
+; CHECK: [[GlobalInvocationId]] = OpVariable [[I32V3PTR]] Input
+; CHECK: [[GlobalSize]] = OpVariable [[I32V3PTR]] Input
+; CHECK: [[GlobalOffset]] = OpVariable [[I32V3PTR]] Input
+; CHECK: [[SubgroupSize]] = OpVariable [[I32PTR]] Input
+; CHECK: [[SubgroupMaxSize]] = OpVariable [[I32PTR]] Input
+; CHECK: [[NumSubgroups]] = OpVariable [[I32PTR]] Input
+; CHECK: [[SubgroupId]] = OpVariable [[I32PTR]] Input
+; CHECK: [[SubgroupLocalInvocationId]] = OpVariable [[I32PTR]] Input
+
+; Function Attrs: convergent noinline norecurse nounwind optnone
+define spir_func void @test_id_and_range() {
+entry:
+  %ssize = alloca i32, align 4
+  %smax = alloca i32, align 4
+  %snum = alloca i32, align 4
+  %sid = alloca i32, align 4
+  %sinvocid = alloca i32, align 4
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[NumWorkgroups]]
+; CHECK: OpCompositeExtract [[I32]] [[LD]] 0
+  %spv.num.workgroups = call i32 @llvm.spv.num.workgroups.i32(i32 0)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[NumWorkgroups]]
+; CHECK: OpCompositeExtract [[I32]] [[LD]] 1
+  %spv.num.workgroups1 = call i32 @llvm.spv.num.workgroups.i32(i32 1)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[NumWorkgroups]]
+; CHECK: OpCompositeExtract [[I32]] [[LD]] 2
+  %spv.num.workgroups2 = call i32 @llvm.spv.num.workgroups.i32(i32 2)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[WorkgroupSize]]
+; CHECK: OpCompositeExtract [[I32]] [[LD]] 0
+  %spv.workgroup.size = call i32 @llvm.spv.workgroup.size.i32(i32 0)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[WorkgroupSize]]
+; CHECK: OpCompositeExtract [[I32]] [[LD]] 1
+  %spv.workgroup.size3 = call i32 @llvm.spv.workgroup.size.i32(i32 1)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[WorkgroupSize]]
+; CHECK: OpCompositeExtract [[I32]] [[LD]] 2
+  %spv.workgroup.size4 = call i32 @llvm.spv.workgroup.size.i32(i32 2)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[WorkgroupId]]
+; CHECK: OpCompositeExtract [[I32]] [[LD]] 0
+  %spv.group.id = call i32 @llvm.spv.group.id.i32(i32 0)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[WorkgroupId]]
+; CHECK: OpCompositeExtract [[I32]] [[LD]] 1
+  %spv.group.id5 = call i32 @llvm.spv.group.id.i32(i32 1)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[WorkgroupId]]
+; CHECK: OpCompositeExtract [[I32]] [[LD]] 2
+  %spv.group.id6 = call i32 @llvm.spv.group.id.i32(i32 2)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[LocalInvocationId]]
+; CHECK: OpCompositeExtract [[I32]] [[LD]] 0
+  %spv.thread.id.in.group = call i32 @llvm.spv.thread.id.in.group.i32(i32 0)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[LocalInvocationId]]
+; CHECK: OpCompositeExtract [[I32]] [[LD]] 1
+  %spv.thread.id.in.group7 = call i32 @llvm.spv.thread.id.in.group.i32(i32 1)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[LocalInvocationId]]
+; CHECK: OpCompositeExtract [[I32]] [[LD]] 2
+  %spv.thread.id.in.group8 = call i32 @llvm.spv.thread.id.in.group.i32(i32 2)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalInvocationId]]
+; CHECK: OpCompositeExtract [[I32]] [[LD]] 0
+  %spv.thread.id = call i32 @llvm.spv.thread.id.i32(i32 0)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalInvocationId]]
+; CHECK: OpCompositeExtract [[I32]] [[LD]] 1
+  %spv.thread.id9 = call i32 @llvm.spv.thread.id.i32(i32 1)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalInvocationId]]
+; CHECK: OpCompositeExtract [[I32]] [[LD]] 2
+  %spv.thread.id10 = call i32 @llvm.spv.thread.id.i32(i32 2)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalSize]]
+; CHECK: OpCompositeExtract [[I32]] [[LD]] 0
+  %spv.num.workgroups11 = call i32 @llvm.spv.global.size.i32(i32 0)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalSize]]
+; CHECK: OpCompositeExtract [[I32]] [[LD]] 1
+  %spv.num.workgroups12 = call i32 @llvm.spv.global.size.i32(i32 1)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalSize]]
+; CHECK: OpCompositeExtract [[I32]] [[LD]] 2
+  %spv.num.workgroups13 = call i32 @llvm.spv.global.size.i32(i32 2)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalOffset]]
+; CHECK: OpCompositeExtract [[I32]] [[LD]] 0
+  %spv.global.offset = call i32 @llvm.spv.global.offset.i32(i32 0)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalOffset]]
+; CHECK: OpCompositeExtract [[I32]] [[LD]] 1
+  %spv.global.offset14 = call i32 @llvm.spv.global.offset.i32(i32 1)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalOffset]]
+; CHECK: OpCompositeExtract [[I32]] [[LD]] 2
+  %spv.global.offset15 = call i32 @llvm.spv.global.offset.i32(i32 2)
+; CHECK: OpLoad %5 [[SubgroupSize]]
+  %0 = call i32 @llvm.spv.subgroup.size()
+  store i32 %0, ptr %ssize, align 4
+; CHECK: OpLoad %5 [[SubgroupMaxSize]]
+  %1 = call i32 @llvm.spv.subgroup.max.size()
+  store i32 %1, ptr %smax, align 4
+; CHECK: OpLoad %5 [[NumSubgroups]]
+  %2 = call i32 @llvm.spv.num.subgroups()
+  store i32 %2, ptr %snum, align 4
+; CHECK: OpLoad %5 [[SubgroupId]]
+  %3 = call i32 @llvm.spv.subgroup.id()
+  store i32 %3, ptr %sid, align 4
+; CHECK: OpLoad %5 [[SubgroupLocalInvocationId]]
+  %4 = call i32 @llvm.spv.subgroup.local.invocation.id()
+  store i32 %4, ptr %sinvocid, align 4
+  ret void
+}
+
+declare i32 @llvm.spv.num.workgroups.i32(i32)
+declare i32 @llvm.spv.workgroup.size.i32(i32)
+declare i32 @llvm.spv.group.id.i32(i32)
+declare i32 @llvm.spv.thread.id.in.group.i32(i32)
+declare i32 @llvm.spv.thread.id.i32(i32)
+declare i32 @llvm.spv.global.size.i32(i32)
+declare i32 @llvm.spv.global.offset.i32(i32)
+declare noundef i32 @llvm.spv.subgroup.size()
+declare noundef i32 @llvm.spv.subgroup.max.size()
+declare noundef i32 @llvm.spv.num.subgroups()
+declare noundef i32 @llvm.spv.subgroup.id()
+declare noundef i32 @llvm.spv.subgroup.local.invocation.id()
diff --git a/llvm/test/CodeGen/SPIRV/builtin_intrinsics_64.ll b/llvm/test/CodeGen/SPIRV/builtin_intrinsics_64.ll
new file mode 100644
index 0000000000000..dcdf8992ce1c4
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/builtin_intrinsics_64.ll
@@ -0,0 +1,137 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1"
+target triple = "spirv64-unknown-unknown"
+
+; CHECK: OpDecorate [[NumWorkgroups:%[0-9]*]] BuiltIn NumWorkgroups
+; CHECK: OpDecorate [[WorkgroupSize:%[0-9]*]] BuiltIn WorkgroupSize
+; CHECK: OpDecorate [[WorkgroupId:%[0-9]*]] BuiltIn WorkgroupId
+; CHECK: OpDecorate [[LocalInvocationId:%[0-9]*]] BuiltIn LocalInvocationId
+; CHECK: OpDecorate [[GlobalInvocationId:%[0-9]*]] BuiltIn GlobalInvocationId
+; CHECK: OpDecorate [[GlobalSize:%[0-9]*]] BuiltIn GlobalSize
+; CHECK: OpDecorate [[GlobalOffset:%[0-9]*]] BuiltIn GlobalOffset
+; CHECK: OpDecorate [[SubgroupSize:%[0-9]*]] BuiltIn SubgroupSize
+; CHECK: OpDecorate [[SubgroupMaxSize:%[0-9]*]] BuiltIn SubgroupMaxSize
+; CHECK: OpDecorate [[NumSubgroups:%[0-9]*]] BuiltIn NumSubgroups
+; CHECK: OpDecorate [[SubgroupId:%[0-9]*]] BuiltIn SubgroupId
+; CHECK: OpDecorate [[SubgroupLocalInvocationId:%[0-9]*]] BuiltIn SubgroupLocalInvocationId
+; CHECK: [[I32:%[0-9]*]] = OpTypeInt 32 0
+; CHECK: [[I64:%[0-9]*]] = OpTypeInt 64 0
+; CHECK: [[I32PTR:%[0-9]*]] = OpTypePointer Input [[I32]]
+; CHECK: [[I64V3:%[0-9]*]] = OpTypeVector [[I64]] 3
+; CHECK: [[I64V3PTR:%[0-9]*]] = OpTypePointer Input [[I64V3]]
+; CHECK: [[NumWorkgroups]] = OpVariable [[I64V3PTR]] Input
+; CHECK: [[WorkgroupSize]] = OpVariable [[I64V3PTR]] Input
+; CHECK: [[WorkgroupId]] = OpVariable [[I64V3PTR]] Input
+; CHECK: [[LocalInvocationId]] = OpVariable [[I64V3PTR]] Input
+; CHECK: [[GlobalInvocationId]] = OpVariable [[I64V3PTR]] Input
+; CHECK: [[GlobalSize]] = OpVariable [[I64V3PTR]] Input
+; CHECK: [[GlobalOffset]] = OpVariable [[I64V3PTR]] Input
+; CHECK: [[SubgroupSize]] = OpVariable [[I32PTR]] Input
+; CHECK: [[SubgroupMaxSize]] = OpVariable [[I32PTR]] Input
+; CHECK: [[NumSubgroups]] = OpVariable [[I32PTR]] Input
+; CHECK: [[SubgroupId]] = OpVariable [[I32PTR]] Input
+; CHECK: [[SubgroupLocalInvocationId]] = OpVariable [[I32PTR]] Input
+
+; Function Attrs: convergent noinline norecurse nounwind optnone
+define spir_func void @test_id_and_range() {
+entry:
+  %ssize = alloca i32, align 4
+  %smax = alloca i32, align 4
+  %snum = alloca i32, align 4
+  %sid = alloca i32, align 4
+  %sinvocid = alloca i32, align 4
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[NumWorkgroups]]
+; CHECK: OpCompositeExtract [[I64]] [[LD]] 0
+  %spv.num.workgroups = call i64 @llvm.spv.num.workgroups.i64(i32 0)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[NumWorkgroups]]
+; CHECK: OpCompositeExtract [[I64]] [[LD]] 1
+  %spv.num.workgroups1 = call i64 @llvm.spv.num.workgroups.i64(i32 1)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[NumWorkgroups]]
+; CHECK: OpCompositeExtract [[I64]] [[LD]] 2
+  %spv.num.workgroups2 = call i64 @llvm.spv.num.workgroups.i64(i32 2)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[WorkgroupSize]]
+; CHECK: OpCompositeExtract [[I64]] [[LD]] 0
+  %spv.workgroup.size = call i64 @llvm.spv.workgroup.size.i64(i32 0)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[WorkgroupSize]]
+; CHECK: OpCompositeExtract [[I64]] [[LD]] 1
+  %spv.workgroup.size3 = call i64 @llvm.spv.workgroup.size.i64(i32 1)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[WorkgroupSize]]
+; CHECK: OpCompositeExtract [[I64]] [[LD]] 2
+  %spv.workgroup.size4 = call i64 @llvm.spv.workgroup.size.i64(i32 2)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[WorkgroupId]]
+; CHECK: OpCompositeExtract [[I64]] [[LD]] 0
+  %spv.group.id = call i64 @llvm.spv.group.id.i64(i32 0)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[WorkgroupId]]
+; CHECK: OpCompositeExtract [[I64]] [[LD]] 1
+  %spv.group.id5 = call i64 @llvm.spv.group.id.i64(i32 1)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[WorkgroupId]]
+; CHECK: OpCompositeExtract [[I64]] [[LD]] 2
+  %spv.group.id6 = call i64 @llvm.spv.group.id.i64(i32 2)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[LocalInvocationId]]
+; CHECK: OpCompositeExtract [[I64]] [[LD]] 0
+  %spv.thread.id.in.group = call i64 @llvm.spv.thread.id.in.group.i64(i32 0)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[LocalInvocationId]]
+; CHECK: OpCompositeExtract [[I64]] [[LD]] 1
+  %spv.thread.id.in.group7 = call i64 @llvm.spv.thread.id.in.group.i64(i32 1)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[LocalInvocationId]]
+; CHECK: OpCompositeExtract [[I64]] [[LD]] 2
+  %spv.thread.id.in.group8 = call i64 @llvm.spv.thread.id.in.group.i64(i32 2)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalInvocationId]]
+; CHECK: OpCompositeExtract [[I64]] [[LD]] 0
+  %spv.thread.id = call i64 @llvm.spv.thread.id.i64(i32 0)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalInvocationId]]
+; CHECK: OpCompositeExtract [[I64]] [[LD]] 1
+  %spv.thread.id9 = call i64 @llvm.spv.thread.id.i64(i32 1)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalInvocationId]]
+; CHECK: OpCompositeExtract [[I64]] [[LD]] 2
+  %spv.thread.id10 = call i64 @llvm.spv.thread.id.i64(i32 2)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalSize]]
+; CHECK: OpCompositeExtract [[I64]] [[LD]] 0
+  %spv.num.workgroups11 = call i64 @llvm.spv.global.size.i64(i32 0)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalSize]]
+; CHECK: OpCompositeExtract [[I64]] [[LD]] 1
+  %spv.num.workgroups12 = call i64 @llvm.spv.global.size.i64(i32 1)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalSize]]
+; CHECK: OpCompositeExtract [[I64]] [[LD]] 2
+  %spv.num.workgroups13 = call i64 @llvm.spv.global.size.i64(i32 2)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalOffset]]
+; CHECK: OpCompositeExtract [[I64]] [[LD]] 0
+  %spv.global.offset = call i64 @llvm.spv.global.offset.i64(i32 0)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalOffset]]
+; CHECK: OpCompositeExtract [[I64]] [[LD]] 1
+  %spv.global.offset14 = call i64 @llvm.spv.global.offset.i64(i32 1)
+; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalOffset]]
+; CHECK: OpCompositeExtract [[I64]] [[LD]] 2
+  %spv.global.offset15 = call i64 @llvm.spv.global.offset.i64(i32 2)
+; CHECK: OpLoad %5 [[SubgroupSize]]
+  %0 = call i32 @llvm.spv.subgroup.size()
+  store i32 %0, ptr %ssize, align 4
+; CHECK: OpLoad %5 [[SubgroupMaxSize]]
+  %1 = call i32 @llvm.spv.subgroup.max.size()
+  store i32 %1, ptr %smax, align 4
+; CHECK: OpLoad %5 [[NumSubgroups]]
+  %2 = call i32 @llvm.spv.num.subgroups()
+  store i32 %2, ptr %snum, align 4
+; CHECK: OpLoad %5 [[SubgroupId]]
+  %3 = call i32 @llvm.spv.subgroup.id()
+  store i32 %3, ptr %sid, align 4
+; CHECK: OpLoad %5 [[SubgroupLocalInvocationId]]
+  %4 = call i32 @llvm.spv.subgroup.local.invocation.id()
+  store i32 %4, ptr %sinvocid, align 4
+  ret void
+}
+
+declare i64 @llvm.spv.num.workgroups.i64(i32)
+declare i64 @llvm.spv.workgroup.size.i64(i32)
+declare i64 @llvm.spv.group.id.i64(i32)
+declare i64 @llvm.spv.thread.id.in.group.i64(i32)
+declare i64 @llvm.spv.thread.id.i64(i32)
+declare i64 @llvm.spv.global.size.i64(i32)
+declare i64 @llvm.spv.global.offset.i64(i32)
+declare noundef i32 @llvm.spv.subgroup.size()
+declare noundef i32 @llvm.spv.subgroup.max.size()
+declare noundef i32 @llvm.spv.num.subgroups()
+declare noundef i32 @llvm.spv.subgroup.id()
+declare noundef i32 @llvm.spv.subgroup.local.invocation.id()
diff --git a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/SV_DispatchThreadID.ll b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/SV_DispatchThreadID.ll
index 2b2ce0974216c..d0d411d2f981d 100644
--- a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/SV_DispatchThreadID.ll
+++ b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/SV_DispatchThreadID.ll
@@ -37,21 +37,21 @@ entry:
 
 ; CHECK:        %[[#load:]] = OpLoad %[[#v3int]] %[[#GlobalInvocationId]]
 ; CHECK:        %[[#load0:]] = OpCompositeExtract %[[#int]] %[[#load]] 0
-  %0 = call i32 @llvm.spv.thread.id(i32 0)
+  %0 = call i32 @llvm.spv.thread.id.i32(i32 0)
 
 ; CHECK:        %[[#tempvar:]] = OpCompositeInsert %[[#v3int]] %[[#load0]] %[[#tempvar]] 0
   %1 = insertelement <3 x i32> poison, i32 %0, i64 0
 
 ; CHECK:        %[[#load:]] = OpLoad %[[#v3int]] %[[#GlobalInvocationId]]
 ; CHECK:        %[[#load1:]] = OpCompositeExtract %[[#int]] %[[#load]] 1
-  %2 = call i32 @llvm.spv.thread.id(i32 1)
+  %2 = call i32 @llvm.spv.thread.id.i32(i32 1)
 
 ; CHECK:        %[[#tempvar:]] = OpCompositeInsert %[[#v3int]] %[[#load1]] %[[#tempvar]] 1
   %3 = insertelement <3 x i32> %1, i32 %2, i64 1
 
 ; CHECK:        %[[#load:]] = OpLoad %[[#v3int]] %[[#GlobalInvocationId]]
 ; CHECK:        %[[#load2:]] = OpCompositeExtract %[[#int]] %[[#load]] 2
-  %4 = call i32 @llvm.spv.thread.id(i32 2)
+  %4 = call i32 @llvm.spv.thread.id.i32(i32 2)
 
 ; CHECK:        %[[#tempvar:]] = OpCompositeInsert %[[#v3int]] %[[#load2]] %[[#tempvar]] 2
   %5 = insertelement <3 x i32> %3, i32 %4, i64 2
@@ -61,7 +61,7 @@ entry:
 }
 
 ; Function Attrs: nounwind willreturn memory(none)
-declare i32 @llvm.spv.thread.id(i32) #2
+declare i32 @llvm.spv.thread.id.i32(i32) #2
 
 attributes #0 = { noinline norecurse nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
 attributes #1 = { norecurse "hlsl.numthreads"="1,1,1" "hlsl.shader"="compute" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
diff --git a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/SV_GroupID.ll b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/SV_GroupID.ll
index bb7650810e989..5b9a7bc02d486 100644
--- a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/SV_GroupID.ll
+++ b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/SV_GroupID.ll
@@ -21,21 +21,21 @@ entry:
 
 ; CHECK:        %[[#load:]] = OpLoad %[[#v3int]] %[[#WorkgroupId]]
 ; CHECK:        %[[#load0:]] = OpCompositeExtract %[[#int]] %[[#load]] 0
-  %1 = call i32 @llvm.spv.group.id(i32 0)
+  %1 = call i32 @llvm.spv.group.id.i32(i32 0)
 
 ; CHECK:        %[[#tempvar:]] = OpCompositeInsert %[[#v3int]] %[[#load0]] %[[#tempvar]]
   %2 = insertelement <3 x i32> poison, i32 %1, i64 0
 
 ; CHECK:        %[[#load:]] = OpLoad %[[#v3int]] %[[#WorkgroupId]]
 ; CHECK:        %[[#load1:]] = OpCompositeExtract %[[#int]] %[[#load]] 1
-  %3 = call i32 @llvm.spv.group.id(i32 1)
+  %3 = call i32 @llvm.spv.group.id.i32(i32 1)
 
 ; CHECK:        %[[#tempvar:]] = OpCompositeInsert %[[#v3int]] %[[#load1]] %[[#tempvar]] 1
   %4 = insertelement <3 x i32> %2, i32 %3, i64 1
 
 ; CHECK:        %[[#load:]] = OpLoad %[[#v3int]] %[[#WorkgroupId]]
 ; CHECK:        %[[#load2:]] = OpCompositeExtract %[[#int]] %[[#load]] 2
-  %5 = call i32 @llvm.spv.group.id(i32 2)
+  %5 = call i32 @llvm.spv.group.id.i32(i32 2)
 
 ; CHECK:        %[[#tempvar:]] = OpCompositeInsert %[[#v3int]] %[[#load2]] %[[#tempvar]] 2
   %6 = insertelement <3 x i32> %4, i32 %5, i64 2
@@ -45,7 +45,7 @@ entry:
 }
 
 ; Function Attrs: nounwind willreturn memory(none)
-declare i32 @llvm.spv.group.id(i32) #3
+declare i32 @llvm.spv.group.id.i32(i32) #3
 
 attributes #1 = { convergent noinline norecurse "hlsl.numthreads"="1,1,1" "hlsl.shader"="compute" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
 attributes #3 = { nounwind willreturn memory(none) }
diff --git a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/SV_GroupThreadID.ll b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/SV_GroupThreadID.ll
index 4e31d3fb77411..f058a539a2263 100644
--- a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/SV_GroupThreadID.ll
+++ b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/SV_GroupThreadID.ll
@@ -37,21 +37,21 @@ entry:
 
 ; CHECK:        %[[#load:]] = OpLoad %[[#v3int]] %[[#LocalInvocationId]]
 ; CHECK:        %[[#load0:]] = OpCompositeExtract %[[#int]] %[[#load]] 0
-  %0 = call i32 @llvm.spv.thread.id.in.group(i32 0)
+  %0 = call i32 @llvm.spv.thread.id.in.group.i32(i32 0)
 
 ; CHECK:        %[[#tempvar:]] = OpCompositeInsert %[[#v3int]] %[[#load0]] %[[#tempvar]] 0
   %1 = insertelement <3 x i32> poison, i32 %0, i64 0
 
 ; CHECK:        %[[#load:]] = OpLoad %[[#v3int]] %[[#LocalInvocationId]]
 ; CHECK:        %[[#load1:]] = OpCompositeExtract %[[#int]] %[[#load]] 1
-  %2 = call i32 @llvm.spv.thread.id.in.group(i32 1)
+  %2 = call i32 @llvm.spv.thread.id.in.group.i32(i32 1)
 
 ; CHECK:        %[[#tempvar:]] = OpCompositeInsert %[[#v3int]] %[[#load1]] %[[#tempvar]] 1
   %3 = insertelement <3 x i32> %1, i32 %2, i64 1
 
 ; CHECK:        %[[#load:]] = OpLoad %[[#v3int]] %[[#LocalInvocationId]]
 ; CHECK:        %[[#load2:]] = OpCompositeExtract %[[#int]] %[[#load]] 2
-  %4 = call i32 @llvm.spv.thread.id.in.group(i32 2)
+  %4 = call i32 @llvm.spv.thread.id.in.group.i32(i32 2)
 
 ; CHECK:        %[[#tempvar:]] = OpCompositeInsert %[[#v3int]] %[[#load2]] %[[#tempvar]] 2
   %5 = insertelement <3 x i32> %3, i32 %4, i64 2
@@ -61,7 +61,7 @@ entry:
 }
 
 ; Function Attrs: nounwind willreturn memory(none)
-declare i32 @llvm.spv.thread.id.in.group(i32) #2
+declare i32 @llvm.spv.thread.id.in.group.i32(i32) #2
 
 attributes #0 = { noinline norecurse nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
 attributes #1 = { norecurse "hlsl.numthreads"="1,1,1" "hlsl.shader"="compute" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }