Skip to content

[SPIRV] Add more id and range builtIns #143909

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 1 commit into
base: main
Choose a base branch
from
Open

[SPIRV] Add more id and range builtIns #143909

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 3 additions & 0 deletions clang/include/clang/Basic/BuiltinsSPIRVCL.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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]>;
10 changes: 10 additions & 0 deletions clang/include/clang/Basic/BuiltinsSPIRVCommon.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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]>;
16 changes: 13 additions & 3 deletions clang/lib/CodeGen/CGHLSLRuntime.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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");
Expand Down
42 changes: 42 additions & 0 deletions clang/lib/CodeGen/TargetBuiltins/SPIR.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
}
40 changes: 39 additions & 1 deletion clang/lib/Headers/__clang_spirv_builtins.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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))
Expand All @@ -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
Expand Down Expand Up @@ -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
Expand Down
8 changes: 5 additions & 3 deletions clang/test/CodeGenHLSL/semantics/DispatchThreadID.hlsl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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]])
Expand Down
18 changes: 12 additions & 6 deletions clang/test/CodeGenHLSL/semantics/SV_GroupID.hlsl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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]])
Expand All @@ -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]])
Expand Down
18 changes: 12 additions & 6 deletions clang/test/CodeGenHLSL/semantics/SV_GroupThreadID.hlsl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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]])
Expand All @@ -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]])
Expand Down
106 changes: 106 additions & 0 deletions clang/test/CodeGenSPIRV/Builtins/ids_and_ranges.c
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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();
}
Loading