llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Victor Lomuller (Naghasan) <details> <summary>Changes</summary> 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. --- 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: - (modified) clang/include/clang/Basic/BuiltinsSPIRVCL.td (+3) - (modified) clang/include/clang/Basic/BuiltinsSPIRVCommon.td (+10) - (modified) clang/lib/CodeGen/CGHLSLRuntime.cpp (+4-3) - (modified) clang/lib/CodeGen/TargetBuiltins/SPIR.cpp (+42) - (modified) clang/lib/Headers/__clang_spirv_builtins.h (+34-1) - (added) clang/test/CodeGenSPIRV/Builtins/ids_and_ranges.c (+106) - (added) clang/test/Headers/spirv_ids.cpp (+110) - (added) clang/test/SemaSPIRV/BuiltIns/ids_and_ranges.c (+77) - (modified) llvm/include/llvm/IR/IntrinsicsSPIRV.td (+18-4) - (modified) llvm/lib/IR/Intrinsics.cpp (+1) - (modified) llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp (+28-2) - (added) llvm/test/CodeGen/SPIRV/builtin_intrinsics_32.ll (+136) - (added) llvm/test/CodeGen/SPIRV/builtin_intrinsics_64.ll (+137) - (modified) llvm/test/CodeGen/SPIRV/hlsl-intrinsics/SV_DispatchThreadID.ll (+4-4) - (modified) llvm/test/CodeGen/SPIRV/hlsl-intrinsics/SV_GroupID.ll (+4-4) - (modified) llvm/test/CodeGen/SPIRV/hlsl-intrinsics/SV_GroupThreadID.ll (+4-4) ``````````diff 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] `````````` </details> https://github.com/llvm/llvm-project/pull/143909 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
