[SPIRV] Add more id and range builtIns

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.

Author: Naghasan

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]>;

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]>;

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");

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;
 }

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

clang/test/CodeGenHLSL/semantics/DispatchThreadID.hlsl

@@ -5,17 +5,19 @@
 
 // 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")]
 [numthreads(8,8,1)]
 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]])

clang/test/CodeGenHLSL/semantics/SV_GroupID.hlsl

@@ -4,17 +4,20 @@
 // 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")]
 [numthreads(8,8,1)]
 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]])

clang/test/CodeGenHLSL/semantics/SV_GroupThreadID.hlsl

@@ -4,17 +4,20 @@
 // 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")]
 [numthreads(8,8,1)]
 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]])

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();
+}