[SPIRV] Add more id and range builtIns #143909
Open
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
llvmbot
added
clang
|
@llvm/pr-subscribers-llvm-ir @llvm/pr-subscribers-clang Author: Victor Lomuller (Naghasan) ChangesThe 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. Patch is 48.65 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/143909.diff 16 Files Affected:
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..ed12a36648367 100644
--- a/clang/lib/CodeGen/CGHLSLRuntime.cpp
+++ b/clang/lib/CodeGen/CGHLSLRuntime.cpp
@@ -394,16 +394,17 @@ llvm::Value *CGHLSLRuntime::emitInputSemantic(IRBuilder<> &B,
}
if (D.hasAttr<HLSLSV_DispatchThreadIDAttr>()) {
llvm::Function *ThreadIDIntrinsic =
- CGM.getIntrinsic(getThreadIdIntrinsic());
+ CGM.getIntrinsic(getThreadIdIntrinsic(), CGM.Int32Ty);
return buildVectorInput(B, ThreadIDIntrinsic, Ty);
}
if (D.hasAttr<HLSLSV_GroupThreadIDAttr>()) {
llvm::Function *GroupThreadIDIntrinsic =
- CGM.getIntrinsic(getGroupThreadIdIntrinsic());
+ CGM.getIntrinsic(getGroupThreadIdIntrinsic(), CGM.Int32Ty);
return buildVectorInput(B, GroupThreadIDIntrinsic, Ty);
}
if (D.hasAttr<HLSLSV_GroupIDAttr>()) {
- llvm::Function *GroupIDIntrinsic = CGM.getIntrinsic(getGroupIdIntrinsic());
+ llvm::Function *GroupIDIntrinsic =
+ CGM.getIntrinsic(getGroupIdIntrinsic(), CGM.Int32Ty);
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..51a0464883b60 100644
--- a/clang/lib/Headers/__clang_spirv_builtins.h
+++ b/clang/lib/Headers/__clang_spirv_builtins.h
@@ -16,6 +16,11 @@
#define __SPIRV_NOEXCEPT
#endif
+#if (!defined(__OPENCL_CPP_VERSION__) && !defined(__OPENCL_C_VERSION__))
+#include <stddef.h>
+#include <stdint.h>
+#endif
+
#define __SPIRV_overloadable __attribute__((overloadable))
#define __SPIRV_convergent __attribute__((convergent))
#define __SPIRV_inline __attribute__((always_inline))
@@ -36,13 +41,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
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...
[truncated]
|
|
Hi, question, why did you prefer passing an new intrinsic to the SPIR-V backend vs loading a global variable with the BuiltIn decoration? For example, when building |
|
The advantage of intrinsics is they can be understood better by passes, it also creates an easier way for passes to introduce them (and less error prone). It's kinda pedantic TBH and I don't have strong opinions about them (it just feels cleaner to me). For the frontend side, the builtin makes things easier and cleaner for SYCL (one of them being that there is no |
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.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
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.