[SPIRV] Add more id and range builtIns (#143909)

The patch adds intrinsics and lowering logic for GlobalSize,
GlobalOffset, SubgroupMaxSize, NumWorkgroups, WorkgroupSize,
WorkgroupId, LocalInvocationId, GlobalInvocationId, SubgroupSize,
NumSubgroups, SubgroupId and SubgroupLocalInvocationId SPIR-V builtins.

The patch also extend spv_thread_id, spv_group_id and
spv_thread_id_in_group to return anyint rather than i32. This allows the
intrinsics to support the opencl environment.

For each of the intrinsics, new clang builtins were added as well as a
binding for the SPIR-V "friendly" format. The original format doesn't
define such binding (uses global variables) but it is not possible to
express the Input SC which is normally required by the environement
specs, and using builtin functions is the most usual approach for other
backend and programming models.

19 files changed, 761 insertions, 37 deletions

diff --git a/clang/include/clang/Basic/BuiltinsSPIRVCL.td b/clang/include/clang/Basic/BuiltinsSPIRVCL.td
index 1103a0d..10320fab3 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 17bcd0b..d2ef6f9 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 c1fd6f2..a47d1cc 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);
   }
   if (D.hasAttr<HLSLSV_PositionAttr>()) {
diff --git a/clang/lib/CodeGen/TargetBuiltins/SPIR.cpp b/clang/lib/CodeGen/TargetBuiltins/SPIR.cpp
index 0687485..1624395 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 e344ed5..9915cdf 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 975a726..7aeb877 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 3aa054a..62985f9 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 3d347b9..2675c97 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 0000000..f71af77
--- /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 0000000..0cd74db
--- /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 0000000..0d98a55
--- /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 43335f8..35c9cd6 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 e631419..d2632d5 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 40a0bd9..fd0bea0 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -3055,6 +3055,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:
@@ -3993,13 +4019,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 0000000..39a755e
--- /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 0000000..dcdf899
--- /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 2b2ce09..d0d411d 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 bb76508..5b9a7bc 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 4e31d3f..f058a53 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" }