https://github.com/Naghasan updated
https://github.com/llvm/llvm-project/pull/143909
>From 049087b07503add00192211b8ee96bdfebcfeed3 Mon Sep 17 00:00:00 2001
From: Victor Lomuller <vic...@codeplay.com>
Date: Thu, 5 Jun 2025 16:17:10 +0100
Subject: [PATCH] [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.
---
clang/include/clang/Basic/BuiltinsSPIRVCL.td | 3 +
.../clang/Basic/BuiltinsSPIRVCommon.td | 10 ++
clang/lib/CodeGen/CGHLSLRuntime.cpp | 16 +-
clang/lib/CodeGen/TargetBuiltins/SPIR.cpp | 42 ++++++
clang/lib/Headers/__clang_spirv_builtins.h | 40 ++++-
.../semantics/DispatchThreadID.hlsl | 8 +-
.../CodeGenHLSL/semantics/SV_GroupID.hlsl | 18 ++-
.../semantics/SV_GroupThreadID.hlsl | 18 ++-
.../CodeGenSPIRV/Builtins/ids_and_ranges.c | 106 ++++++++++++++
clang/test/Headers/spirv_ids.cpp | 110 ++++++++++++++
.../test/SemaSPIRV/BuiltIns/ids_and_ranges.c | 77 ++++++++++
llvm/include/llvm/IR/IntrinsicsSPIRV.td | 22 ++-
llvm/lib/IR/Intrinsics.cpp | 1 +
.../Target/SPIRV/SPIRVInstructionSelector.cpp | 30 +++-
.../CodeGen/SPIRV/builtin_intrinsics_32.ll | 136 +++++++++++++++++
.../CodeGen/SPIRV/builtin_intrinsics_64.ll | 137 ++++++++++++++++++
.../hlsl-intrinsics/SV_DispatchThreadID.ll | 8 +-
.../SPIRV/hlsl-intrinsics/SV_GroupID.ll | 8 +-
.../SPIRV/hlsl-intrinsics/SV_GroupThreadID.ll | 8 +-
19 files changed, 761 insertions(+), 37 deletions(-)
create mode 100644 clang/test/CodeGenSPIRV/Builtins/ids_and_ranges.c
create mode 100644 clang/test/Headers/spirv_ids.cpp
create mode 100644 clang/test/SemaSPIRV/BuiltIns/ids_and_ranges.c
create mode 100644 llvm/test/CodeGen/SPIRV/builtin_intrinsics_32.ll
create mode 100644 llvm/test/CodeGen/SPIRV/builtin_intrinsics_64.ll
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" }
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
