aboutsummaryrefslogtreecommitdiff
path: root/clang
diff options
context:
space:
mode:
Diffstat (limited to 'clang')
-rw-r--r--clang/include/clang/Basic/BuiltinsSPIRVCL.td3
-rw-r--r--clang/include/clang/Basic/BuiltinsSPIRVCommon.td10
-rw-r--r--clang/lib/CodeGen/CGHLSLRuntime.cpp16
-rw-r--r--clang/lib/CodeGen/TargetBuiltins/SPIR.cpp42
-rw-r--r--clang/lib/Headers/__clang_spirv_builtins.h40
-rw-r--r--clang/test/CodeGenHLSL/semantics/DispatchThreadID.hlsl8
-rw-r--r--clang/test/CodeGenHLSL/semantics/SV_GroupID.hlsl18
-rw-r--r--clang/test/CodeGenHLSL/semantics/SV_GroupThreadID.hlsl18
-rw-r--r--clang/test/CodeGenSPIRV/Builtins/ids_and_ranges.c106
-rw-r--r--clang/test/Headers/spirv_ids.cpp110
-rw-r--r--clang/test/SemaSPIRV/BuiltIns/ids_and_ranges.c77
11 files changed, 429 insertions, 19 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}}
+}
generated by cgit v1.1 at 2025-08-04 14:37:34 +0000