[llvm] cda4820 - [SPIRV] Do not use OpTypeRuntimeArray in Kernel env. (#149522)
via llvm-commits
llvm-commits at lists.llvm.org
Tue Aug 5 06:10:19 PDT 2025
Author: Marcos Maronas
Date: 2025-08-05T15:10:15+02:00
New Revision: cda4820270dc9d4e84107152130a54c5ecc4976d
URL: https://github.com/llvm/llvm-project/commit/cda4820270dc9d4e84107152130a54c5ecc4976d
DIFF: https://github.com/llvm/llvm-project/commit/cda4820270dc9d4e84107152130a54c5ecc4976d.diff
LOG: [SPIRV] Do not use OpTypeRuntimeArray in Kernel env. (#149522)
Prior to this patch, when `NumElems` was 0, `OpTypeRuntimeArray` was
directly generated, but it requires `Shader` capability, so it can only
be generated if `Shader` env is being used. We have observed a pattern
of using unbound arrays that translate into `[0 x ...]` types in OpenCL,
which implies `Kernel` capability, so `OpTypeRuntimeArray` should not be
used. To prevent this scenario, this patch simplifies GEP instructions
where type is a 0-length array and the first index is also 0. In such
scenario, we effectively drop the 0-length array and the first index.
Additionally, the newly added test prior to this patch was generating a
module with both `Shader` and `Kernel` capabilities at the same time,
but they're incompatible. This patch also fixes that.
Finally, prior to this patch, the newly added test was adding `Shader`
capability to the module even with the command line flag
`--avoid-spirv-capabilities=Shader`. This patch also has a fix for that.
Added:
llvm/test/CodeGen/SPIRV/array_type.ll
llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions/cl_khr_extended_bit_ops_spv-friendly_only_no_extension.ll
llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions_no_extension.ll
llvm/test/CodeGen/SPIRV/pointers/variables-storage-class-private.ll
llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_v2i32.ll
Modified:
llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions.ll
llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions/cl_khr_extended_bit_ops_spv-friendly_only.ll
llvm/test/CodeGen/SPIRV/llvm-intrinsics/bitreverse_small_type.ll
llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll
llvm/test/CodeGen/SPIRV/memory_model_md.ll
llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.ll
llvm/test/CodeGen/SPIRV/pointers/variables-storage-class.ll
llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_i32.ll
llvm/test/CodeGen/SPIRV/zero-length-array.ll
Removed:
llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_v2i16.ll
################################################################################
diff --git a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
index 9f55330cb7ac6..2c3e0876b757d 100644
--- a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
@@ -21,7 +21,9 @@
#include "llvm/IR/InstIterator.h"
#include "llvm/IR/InstVisitor.h"
#include "llvm/IR/IntrinsicsSPIRV.h"
+#include "llvm/IR/PatternMatch.h"
#include "llvm/IR/TypedPointerType.h"
+#include "llvm/Transforms/Utils/Local.h"
#include <queue>
#include <unordered_set>
@@ -187,6 +189,8 @@ class SPIRVEmitIntrinsics
void applyDemangledPtrArgTypes(IRBuilder<> &B);
+ GetElementPtrInst *simplifyZeroLengthArrayGepInst(GetElementPtrInst *GEP);
+
bool runOnFunction(Function &F);
bool postprocessTypes(Module &M);
bool processFunctionPointers(Module &M);
@@ -2561,6 +2565,30 @@ void SPIRVEmitIntrinsics::applyDemangledPtrArgTypes(IRBuilder<> &B) {
}
}
+GetElementPtrInst *
+SPIRVEmitIntrinsics::simplifyZeroLengthArrayGepInst(GetElementPtrInst *GEP) {
+ // getelementptr [0 x T], P, 0 (zero), I -> getelementptr T, P, I.
+ // If type is 0-length array and first index is 0 (zero), drop both the
+ // 0-length array type and the first index. This is a common pattern in the
+ // IR, e.g. when using a zero-length array as a placeholder for a flexible
+ // array such as unbound arrays.
+ assert(GEP && "GEP is null");
+ Type *SrcTy = GEP->getSourceElementType();
+ SmallVector<Value *, 8> Indices(GEP->indices());
+ ArrayType *ArrTy = dyn_cast<ArrayType>(SrcTy);
+ if (ArrTy && ArrTy->getNumElements() == 0 &&
+ PatternMatch::match(Indices[0], PatternMatch::m_Zero())) {
+ IRBuilder<> Builder(GEP);
+ Indices.erase(Indices.begin());
+ SrcTy = ArrTy->getElementType();
+ Value *NewGEP = Builder.CreateGEP(SrcTy, GEP->getPointerOperand(), Indices,
+ "", GEP->getNoWrapFlags());
+ assert(llvm::isa<GetElementPtrInst>(NewGEP) && "NewGEP should be a GEP");
+ return cast<GetElementPtrInst>(NewGEP);
+ }
+ return nullptr;
+}
+
bool SPIRVEmitIntrinsics::runOnFunction(Function &Func) {
if (Func.isDeclaration())
return false;
@@ -2578,14 +2606,30 @@ bool SPIRVEmitIntrinsics::runOnFunction(Function &Func) {
AggrConstTypes.clear();
AggrStores.clear();
- // fix GEP result types ahead of inference
+ // Fix GEP result types ahead of inference, and simplify if possible.
+ // Data structure for dead instructions that were simplified and replaced.
+ SmallPtrSet<Instruction *, 4> DeadInsts;
for (auto &I : instructions(Func)) {
auto *Ref = dyn_cast<GetElementPtrInst>(&I);
if (!Ref || GR->findDeducedElementType(Ref))
continue;
+
+ GetElementPtrInst *NewGEP = simplifyZeroLengthArrayGepInst(Ref);
+ if (NewGEP) {
+ Ref->replaceAllUsesWith(NewGEP);
+ if (isInstructionTriviallyDead(Ref))
+ DeadInsts.insert(Ref);
+
+ Ref = NewGEP;
+ }
if (Type *GepTy = getGEPType(Ref))
GR->addDeducedElementType(Ref, normalizeType(GepTy));
}
+ // Remove dead instructions that were simplified and replaced.
+ for (auto *I : DeadInsts) {
+ assert(I->use_empty() && "Dead instruction should not have any uses left");
+ I->eraseFromParent();
+ }
processParamTypesByFunHeader(CurrF, B);
diff --git a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
index 83fccdc2bdba3..960eb2ef93a9e 100644
--- a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
@@ -828,6 +828,8 @@ SPIRVType *SPIRVGlobalRegistry::getOpTypeArray(uint32_t NumElems,
"Invalid array element type");
SPIRVType *SpvTypeInt32 = getOrCreateSPIRVIntegerType(32, MIRBuilder);
SPIRVType *ArrayType = nullptr;
+ const SPIRVSubtarget &ST =
+ cast<SPIRVSubtarget>(MIRBuilder.getMF().getSubtarget());
if (NumElems != 0) {
Register NumElementsVReg =
buildConstantInt(NumElems, MIRBuilder, SpvTypeInt32, EmitIR);
@@ -838,6 +840,8 @@ SPIRVType *SPIRVGlobalRegistry::getOpTypeArray(uint32_t NumElems,
.addUse(NumElementsVReg);
});
} else {
+ if (!ST.isShader())
+ return nullptr;
ArrayType = createOpType(MIRBuilder, [&](MachineIRBuilder &MIRBuilder) {
return MIRBuilder.buildInstr(SPIRV::OpTypeRuntimeArray)
.addDef(createTypeVReg(MIRBuilder))
diff --git a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
index 0cd9d7882a52a..000e2d7105c8d 100644
--- a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
@@ -744,8 +744,14 @@ void SPIRV::RequirementHandler::checkSatisfiable(
IsSatisfiable = false;
}
+ AvoidCapabilitiesSet AvoidCaps;
+ if (!ST.isShader())
+ AvoidCaps.S.insert(SPIRV::Capability::Shader);
+ else
+ AvoidCaps.S.insert(SPIRV::Capability::Kernel);
+
for (auto Cap : MinimalCaps) {
- if (AvailableCaps.contains(Cap))
+ if (AvailableCaps.contains(Cap) && !AvoidCaps.S.contains(Cap))
continue;
LLVM_DEBUG(dbgs() << "Capability not supported: "
<< getSymbolicOperandMnemonic(
@@ -1865,6 +1871,10 @@ void addInstrRequirements(const MachineInstr &MI,
Reqs.addCapability(SPIRV::Capability::TernaryBitwiseFunctionINTEL);
break;
}
+ case SPIRV::OpCopyMemorySized: {
+ Reqs.addCapability(SPIRV::Capability::Addresses);
+ // TODO: Add UntypedPointersKHR when implemented.
+ }
default:
break;
diff --git a/llvm/test/CodeGen/SPIRV/array_type.ll b/llvm/test/CodeGen/SPIRV/array_type.ll
new file mode 100644
index 0000000000000..251b48f8bf629
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/array_type.ll
@@ -0,0 +1,78 @@
+; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64v1.2-unknown-unknown %s -o - -filetype=obj | spirv-val --target-env opencl2.2 %}
+
+; CHECK: OpCapability Kernel
+; CHECK-NOT: OpCapability Shader
+; CHECK-DAG: %[[#float16:]] = OpTypeFloat 16
+; CHECK-DAG: %[[#SyclHalfTy:]] = OpTypeStruct %[[#float16]]
+; CHECK-DAG: %[[#i16:]] = OpTypeInt 16
+; CHECK-DAG: %[[#i32:]] = OpTypeInt 32
+; CHECK-DAG: %[[#i64:]] = OpTypeInt 64
+; CHECK-DAG: %[[#ConstNull:]] = OpConstantNull %[[#i64]]
+; CHECK-DAG: %[[#ConstOne:]] = OpConstant %[[#i64]] 1
+; CHECK-DAG: %[[#ConstFive:]] = OpConstant %[[#i16]] 5
+; CHECK-DAG: %[[#SyclHalfTyPtr:]] = OpTypePointer Function %[[#SyclHalfTy]]
+; CHECK-DAG: %[[#i32Ptr:]] = OpTypePointer Function %[[#i32]]
+; CHECK-DAG: %[[#StorePtrTy:]] = OpTypePointer Function %[[#i16]]
+
+%"class.sycl::_V1::detail::half_impl::half" = type { half }
+
+; Function Attrs: mustprogress norecurse nounwind
+define spir_kernel void @foo(ptr %p){
+; CHECK: OpFunction
+; CHECK: %[[#Ptr:]] = OpFunctionParameter
+; CHECK: OpLabel
+; CHECK: %[[#BitcastOp:]] = OpInBoundsPtrAccessChain %[[#SyclHalfTyPtr]] %[[#Ptr]] %[[#ConstNull]] %[[#ConstNull]]
+; CHECK: %[[#StorePtr:]] = OpBitcast %[[#StorePtrTy]] %[[#BitcastOp]]
+; CHECK: OpStore %[[#StorePtr]] %[[#ConstFive]]
+; CHECK: OpReturn
+entry:
+ %0 = getelementptr inbounds [0 x [32 x %"class.sycl::_V1::detail::half_impl::half"]], ptr %p, i64 0, i64 0, i64 0
+ store i16 5, ptr %0
+ ret void
+}
+
+; Function Attrs: mustprogress norecurse nounwind
+define spir_kernel void @foo2(ptr %p){
+; CHECK: OpFunction
+; CHECK: %[[#Ptr:]] = OpFunctionParameter
+; CHECK: OpLabel
+; CHECK: %[[#BitcastOp:]] = OpInBoundsPtrAccessChain %[[#SyclHalfTyPtr]] %[[#Ptr]] %[[#ConstOne]] %[[#ConstOne]]
+; CHECK: %[[#StorePtr:]] = OpBitcast %[[#StorePtrTy]] %[[#BitcastOp]]
+; CHECK: OpStore %[[#StorePtr]] %[[#ConstFive]]
+; CHECK: OpReturn
+entry:
+ %0 = getelementptr inbounds [0 x [32 x %"class.sycl::_V1::detail::half_impl::half"]], ptr %p, i64 0, i64 1, i64 1
+ store i16 5, ptr %0
+ ret void
+}
+
+; Function Attrs: mustprogress norecurse nounwind
+define spir_kernel void @foo3(ptr %p){
+; CHECK: OpFunction
+; CHECK: %[[#Ptr:]] = OpFunctionParameter
+; CHECK: OpLabel
+; CHECK: %[[#BitcastOp:]] = OpInBoundsPtrAccessChain %[[#i32Ptr]] %[[#Ptr]] %[[#ConstNull]] %[[#ConstNull]]
+; CHECK: %[[#StorePtr:]] = OpBitcast %[[#StorePtrTy]] %[[#BitcastOp]]
+; CHECK: OpStore %[[#StorePtr]] %[[#ConstFive]]
+; CHECK: OpReturn
+entry:
+ %0 = getelementptr inbounds [0 x [32 x i32]], ptr %p, i64 0, i64 0, i64 0
+ store i16 5, ptr %0
+ ret void
+}
+
+; Function Attrs: mustprogress norecurse nounwind
+define spir_kernel void @foo4(ptr %p){
+; CHECK: OpFunction
+; CHECK: %[[#Ptr:]] = OpFunctionParameter
+; CHECK: OpLabel
+; CHECK: %[[#BitcastOp:]] = OpInBoundsPtrAccessChain %[[#i32Ptr]] %[[#Ptr]] %[[#ConstOne]] %[[#ConstOne]]
+; CHECK: %[[#StorePtr:]] = OpBitcast %[[#StorePtrTy]] %[[#BitcastOp]]
+; CHECK: OpStore %[[#StorePtr]] %[[#ConstFive]]
+; CHECK: OpReturn
+entry:
+ %0 = getelementptr inbounds [0 x [32 x i32]], ptr %p, i64 0, i64 1, i64 1
+ store i16 5, ptr %0
+ ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions.ll
index 40e2aff0d755a..7adb039464c4f 100644
--- a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions.ll
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions.ll
@@ -1,18 +1,11 @@
; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s --spirv-ext=+SPV_KHR_bit_instructions -o - | FileCheck %s --check-prefix=CHECK-EXTENSION
-; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-NO-EXTENSION
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64v1.2-unknown-unknown %s --spirv-ext=+SPV_KHR_bit_instructions -o - -filetype=obj | spirv-val --target-env opencl2.2 %}
; CHECK-EXTENSION: OpCapability BitInstructions
; CHECK-EXTENSION-NEXT: OpExtension "SPV_KHR_bit_instructions"
; CHECK-EXTENSION-NOT: OpCabilitity Shader
-; CHECK-NO-EXTENSION: OpCapability Shader
-; CHECK-NO-EXTENSION-NOT: OpCabilitity BitInstructions
-; CHECK-NO-EXTENSION-NOT: OpExtension "SPV_KHR_bit_instructions"
-
-
; CHECK-EXTENSION: %[[#int:]] = OpTypeInt 32
; CHECK-EXTENSION: OpBitReverse %[[#int]]
-; CHECK-NO-EXTENSION: %[[#int:]] = OpTypeInt 32
-; CHECK-NO-EXTENSION: OpBitReverse %[[#int]]
define spir_kernel void @testBitRev(i32 %a, i32 %b, i32 %c, i32 addrspace(1)* nocapture %res) local_unnamed_addr {
entry:
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions/cl_khr_extended_bit_ops_spv-friendly_only.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions/cl_khr_extended_bit_ops_spv-friendly_only.ll
index 65cccc83a3e02..3bd1bd633c258 100644
--- a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions/cl_khr_extended_bit_ops_spv-friendly_only.ll
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions/cl_khr_extended_bit_ops_spv-friendly_only.ll
@@ -1,12 +1,8 @@
; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s --spirv-ext=+SPV_KHR_bit_instructions -o - | FileCheck %s --check-prefix=CHECK-EXTENSION
-; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-NO-EXTENSION
-; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s --spirv-ext=+SPV_KHR_bit_instructions -o - -filetype=obj | spirv-val %}
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64v1.2-unknown-unknown %s --spirv-ext=+SPV_KHR_bit_instructions -o - -filetype=obj | spirv-val --target-env opencl2.2 %}
;
; CHECK-EXTENSION: Capability BitInstructions
; CHECK-EXTENSION: Extension "SPV_KHR_bit_instructions"
-; CHECK-NO-EXTENSION-NOT: Capability BitInstructions
-; CHECK-NO-EXTENSION-NOT: Extension "SPV_KHR_bit_instructions"
-; CHECK-NO-EXTENSION: Capability Shader
;
; CHECK-EXTENSION: %[[#]] = OpFunction %[[#]] None %[[#]]
; CHECK-EXTENSION: %[[#reversebase:]] = OpFunctionParameter %[[#]]
@@ -15,24 +11,11 @@
; kernel void testBitReverse_SPIRVFriendly(long4 b, global long4 *res) {
; *res = bit_reverse(b);
; }
-define spir_kernel void @testBitReverse_SPIRVFriendly(<4 x i64> %b, ptr addrspace(1) nocapture align 32 %res) #3 {
+define spir_kernel void @testBitReverse_SPIRVFriendly(<4 x i64> %b, ptr addrspace(1) %res) {
entry:
%call = call <4 x i64> @llvm.bitreverse.v4i64(<4 x i64> %b)
- store <4 x i64> %call, ptr addrspace(1) %res, align 32
+ store <4 x i64> %call, ptr addrspace(1) %res
ret void
}
-declare <4 x i64> @llvm.bitreverse.v4i64(<4 x i64>) #4
-
-
-attributes #3 = { nounwind }
-attributes #4 = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
-
-!llvm.module.flags = !{!0}
-!opencl.ocl.version = !{!1}
-!opencl.spir.version = !{!1}
-!llvm.ident = !{!2}
-
-!0 = !{i32 1, !"wchar_size", i32 4}
-!1 = !{i32 2, i32 0}
-!2 = !{!"clang version 20.0.0git (https://github.com/llvm/llvm-project.git cc61409d353a40f62d3a137f3c7436aa00df779d)"}
+declare <4 x i64> @llvm.bitreverse.v4i64(<4 x i64>)
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions/cl_khr_extended_bit_ops_spv-friendly_only_no_extension.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions/cl_khr_extended_bit_ops_spv-friendly_only_no_extension.ll
new file mode 100644
index 0000000000000..61ef273a25734
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions/cl_khr_extended_bit_ops_spv-friendly_only_no_extension.ll
@@ -0,0 +1,16 @@
+; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-NO-EXTENSION
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+;
+; CHECK-NO-EXTENSION-NOT: Capability BitInstructions
+; CHECK-NO-EXTENSION-NOT: Extension "SPV_KHR_bit_instructions"
+; CHECK-NO-EXTENSION: Capability Shader
+
+define internal spir_func void @testBitReverse_SPIRVFriendly() #3 {
+entry:
+ %call = call <4 x i64> @llvm.bitreverse.v4i64(<4 x i64> <i64 1, i64 2, i64 3, i64 4>)
+ ret void
+}
+
+declare <4 x i64> @llvm.bitreverse.v4i64(<4 x i64>)
+
+attributes #3 = { nounwind "hlsl.shader"="compute" }
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions_no_extension.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions_no_extension.ll
new file mode 100644
index 0000000000000..452df0a64c063
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions_no_extension.ll
@@ -0,0 +1,23 @@
+; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - | FileCheck %s --check-prefix=CHECK-NO-EXTENSION
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - -filetype=obj | spirv-val --target-env vulkan1.3 %}
+
+
+; CHECK-NO-EXTENSION: OpCapability Shader
+; CHECK-NO-EXTENSION-NOT: OpCabilitity BitInstructions
+; CHECK-NO-EXTENSION-NOT: OpExtension "SPV_KHR_bit_instructions"
+; CHECK-NO-EXTENSION: %[[#int:]] = OpTypeInt 32
+; CHECK-NO-EXTENSION: OpBitReverse %[[#int]]
+
+define hidden spir_func void @testBitRev(i32 %a, i32 %b, i32 %c, ptr %res) local_unnamed_addr {
+entry:
+ %call = tail call i32 @llvm.bitreverse.i32(i32 %b)
+ store i32 %call, ptr %res, align 4
+ ret void
+}
+
+define void @main() #1 {
+ ret void
+}
+
+declare i32 @llvm.bitreverse.i32(i32)
+attributes #1 = { "hlsl.numthreads"="8,1,1" "hlsl.shader"="compute" }
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/bitreverse_small_type.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/bitreverse_small_type.ll
index 77b8c5118e59b..438fff6e94f89 100644
--- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/bitreverse_small_type.ll
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/bitreverse_small_type.ll
@@ -1,8 +1,8 @@
;; Check that llvm.bitreverse.* intrinsics are lowered for
;; 2/4-bit scalar and vector types.
-; RUN: llc -O0 -verify-machineinstrs -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_INTEL_arbitrary_precision_integers %s -o - | FileCheck %s
-; TODO: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_INTEL_arbitrary_precision_integers %s -o - -filetype=obj | spirv-val %}
+; RUN: llc -O0 -verify-machineinstrs -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_INTEL_arbitrary_precision_integers,+SPV_KHR_bit_instructions %s -o - | FileCheck %s
+; TODO: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_INTEL_arbitrary_precision_integers,+SPV_KHR_bit_instructions %s -o - -filetype=obj | spirv-val %}
; CHECK: OpCapability ArbitraryPrecisionIntegersINTEL
; CHECK: OpExtension "SPV_INTEL_arbitrary_precision_integers"
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll
index 9d07b63b49a52..483d7077a66f9 100644
--- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll
@@ -18,7 +18,6 @@
; CL: %[[#FooVar:]] = OpVariable
; CL-NEXT: %[[#Casted1:]] = OpBitcast %[[#PtrChar]] %[[#FooVar]]
; CL-NEXT: OpLifetimeStart %[[#Casted1]], 72
-; CL-NEXT: OpCopyMemorySized
; CL-NEXT: OpBitcast
; CL-NEXT: OpInBoundsPtrAccessChain
; CL-NEXT: %[[#Casted2:]] = OpBitcast %[[#PtrChar]] %[[#FooVar]]
@@ -26,13 +25,11 @@
; VK: OpFunction
; VK: %[[#FooVar:]] = OpVariable
-; VK-NEXT: OpCopyMemorySized
; VK-NEXT: OpInBoundsAccessChain
; VK-NEXT: OpReturn
define spir_func void @foo(ptr noundef byval(%tprange) align 8 %_arg_UserRange) {
%RoundedRangeKernel = alloca %tprange, align 8
call void @llvm.lifetime.start.p0(i64 72, ptr nonnull %RoundedRangeKernel)
- call void @llvm.memcpy.p0.p0.i64(ptr align 8 %RoundedRangeKernel, ptr align 8 %_arg_UserRange, i64 16, i1 false)
%KernelFunc = getelementptr inbounds i8, ptr %RoundedRangeKernel, i64 8
call void @llvm.lifetime.end.p0(i64 72, ptr nonnull %RoundedRangeKernel)
ret void
@@ -41,20 +38,17 @@ define spir_func void @foo(ptr noundef byval(%tprange) align 8 %_arg_UserRange)
; CL: OpFunction
; CL: %[[#BarVar:]] = OpVariable
; CL-NEXT: OpLifetimeStart %[[#BarVar]], 0
-; CL-NEXT: OpCopyMemorySized
; CL-NEXT: OpBitcast
; CL-NEXT: OpInBoundsPtrAccessChain
; CL-NEXT: OpLifetimeStop %[[#BarVar]], 0
; VK: OpFunction
; VK: %[[#BarVar:]] = OpVariable
-; VK-NEXT: OpCopyMemorySized
; VK-NEXT: OpInBoundsAccessChain
; VK-NEXT: OpReturn
define spir_func void @bar(ptr noundef byval(%tprange) align 8 %_arg_UserRange) {
%RoundedRangeKernel = alloca %tprange, align 8
call void @llvm.lifetime.start.p0(i64 -1, ptr nonnull %RoundedRangeKernel)
- call void @llvm.memcpy.p0.p0.i64(ptr align 8 %RoundedRangeKernel, ptr align 8 %_arg_UserRange, i64 16, i1 false)
%KernelFunc = getelementptr inbounds i8, ptr %RoundedRangeKernel, i64 8
call void @llvm.lifetime.end.p0(i64 -1, ptr nonnull %RoundedRangeKernel)
ret void
@@ -63,20 +57,17 @@ define spir_func void @bar(ptr noundef byval(%tprange) align 8 %_arg_UserRange)
; CL: OpFunction
; CL: %[[#TestVar:]] = OpVariable
; CL-NEXT: OpLifetimeStart %[[#TestVar]], 1
-; CL-NEXT: OpCopyMemorySized
; CL-NEXT: OpInBoundsPtrAccessChain
; CL-NEXT: OpLifetimeStop %[[#TestVar]], 1
; VK: OpFunction
; VK: %[[#Test:]] = OpVariable
-; VK-NEXT: OpCopyMemorySized
; VK-NEXT: OpInBoundsAccessChain
; VK-NEXT: OpReturn
define spir_func void @test(ptr noundef align 8 %_arg) {
%var = alloca i8, align 8
call void @llvm.lifetime.start.p0(i64 1, ptr nonnull %var)
- call void @llvm.memcpy.p0.p0.i64(ptr align 8 %var, ptr align 8 %_arg, i64 1, i1 false)
- %KernelFunc = getelementptr inbounds i8, ptr %var, i64 0
+ %KernelFunc = getelementptr inbounds i8, ptr %var, i64 1
call void @llvm.lifetime.end.p0(i64 1, ptr nonnull %var)
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/memory_model_md.ll b/llvm/test/CodeGen/SPIRV/memory_model_md.ll
index e52343cbbb7e4..684a163397ca8 100644
--- a/llvm/test/CodeGen/SPIRV/memory_model_md.ll
+++ b/llvm/test/CodeGen/SPIRV/memory_model_md.ll
@@ -1,6 +1,7 @@
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPV
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32v1.2-unknown-unknown %s -o - -filetype=obj | spirv-val --target-env opencl2.2 %}
-; SPV: OpMemoryModel Physical32 Simple
+; SPV: OpMemoryModel Physical32 OpenCL
define dso_local dllexport void @k_no_fc(i32 %ibuf, i32 %obuf) local_unnamed_addr {
entry:
ret void
@@ -8,4 +9,4 @@ entry:
!spirv.MemoryModel = !{!0}
-!0 = !{i32 1, i32 0}
+!0 = !{i32 1, i32 2}
diff --git a/llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.ll b/llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.ll
index 544c657da8488..19451d23c6830 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.ll
@@ -1,5 +1,5 @@
-; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
-; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - -filetype=obj | spirv-val --target-env vulkan1.3 %}
@PrivInternal = internal addrspace(10) global i32 456
; CHECK-DAG: %[[#type:]] = OpTypeInt 32 0
@@ -7,7 +7,7 @@
; CHECK-DAG: %[[#value:]] = OpConstant %[[#type]] 456
; CHECK-DAG: %[[#var:]] = OpVariable %[[#ptrty]] Private %[[#value]]
-define spir_kernel void @Foo() {
+define hidden spir_func void @Foo() {
%p = addrspacecast ptr addrspace(10) @PrivInternal to ptr
%v = load i32, ptr %p, align 4
ret void
@@ -15,3 +15,9 @@ define spir_kernel void @Foo() {
; CHECK-NEXT: OpLoad %[[#type]] %[[#var]] Aligned 4
; CHECK-Next: OpReturn
}
+
+define void @main() #1 {
+ ret void
+}
+
+attributes #1 = { "hlsl.numthreads"="8,1,1" "hlsl.shader"="compute" }
diff --git a/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class-private.ll b/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class-private.ll
new file mode 100644
index 0000000000000..51db12046e9de
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class-private.ll
@@ -0,0 +1,21 @@
+; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - -filetype=obj | spirv-val --target-env vulkan1.3%}
+
+; CHECK-DAG: %[[#U32:]] = OpTypeInt 32 0
+
+; CHECK-DAG: %[[#VAL:]] = OpConstant %[[#U32]] 456
+; CHECK-DAG: %[[#VTYPE:]] = OpTypePointer Private %[[#U32]]
+; CHECK-DAG: %[[#]] = OpVariable %[[#VTYPE]] Private %[[#VAL]]
+ at PrivInternal = internal addrspace(10) global i32 456
+
+define hidden spir_func void @Foo() {
+ %tmp = load i32, ptr addrspace(10) @PrivInternal
+ ret void
+}
+
+define void @main() #1 {
+ ret void
+}
+
+declare void @llvm.memcpy.p1.p2.i64(ptr addrspace(1) noalias nocapture writeonly, ptr addrspace(2) noalias nocapture readonly, i64, i1 immarg)
+attributes #1 = { "hlsl.numthreads"="8,1,1" "hlsl.shader"="compute" }
diff --git a/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class.ll b/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class.ll
index a1ded0569d67e..6914f4faebdcd 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class.ll
@@ -1,5 +1,5 @@
; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
-; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64v1.2-unknown-unknown %s -o - -filetype=obj | spirv-val --target-env opencl2.2 %}
; CHECK-DAG: %[[#U8:]] = OpTypeInt 8 0
; CHECK-DAG: %[[#U32:]] = OpTypeInt 32 0
@@ -15,12 +15,7 @@
; CHECK-DAG: %[[#INIT:]] = OpVariable %[[#VTYPE]] UniformConstant %[[#VAL]]
@Init = private addrspace(2) constant i32 123
-; CHECK-DAG: %[[#VAL:]] = OpConstant %[[#U32]] 456
-; CHECK-DAG: %[[#VTYPE:]] = OpTypePointer Private %[[#U32]]
-; CHECK-DAG: %[[#]] = OpVariable %[[#VTYPE]] Private %[[#VAL]]
- at PrivInternal = internal addrspace(10) global i32 456
-
-define spir_kernel void @Foo() {
+define internal spir_func void @Foo() {
; CHECK: %[[#]] = OpLoad %[[#]] %[[#PTR]] Aligned 8
%l = load ptr addrspace(1), ptr addrspace(1) @Ptr, align 8
; CHECK: OpCopyMemorySized %[[#]] %[[#INIT]] %[[#]] Aligned 4
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_i32.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_i32.ll
index f396b5a01ae91..838c55172d3fc 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_i32.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_i32.ll
@@ -1,13 +1,19 @@
-; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: llc -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - -filetype=obj | spirv-val --target-env vulkan1.3 %}
; CHECK-SPIRV: %[[#int:]] = OpTypeInt 32
; CHECK-SPIRV: OpBitReverse %[[#int]]
-define spir_kernel void @testBitRev(i32 %a, i32 %b, i32 %c, i32 addrspace(1)* nocapture %res) local_unnamed_addr {
+define hidden spir_func void @testBitRev(i32 %a, i32 %b, i32 %c, ptr %res) {
entry:
%call = tail call i32 @llvm.bitreverse.i32(i32 %b)
- store i32 %call, i32 addrspace(1)* %res, align 4
+ store i32 %call, ptr %res, align 4
+ ret void
+}
+
+define void @main() #1 {
ret void
}
declare i32 @llvm.bitreverse.i32(i32)
+attributes #1 = { "hlsl.numthreads"="8,1,1" "hlsl.shader"="compute" }
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_v2i16.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_v2i16.ll
deleted file mode 100644
index 8f04929fdd587..0000000000000
--- a/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_v2i16.ll
+++ /dev/null
@@ -1,14 +0,0 @@
-; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-
-; CHECK-SPIRV: %[[#short:]] = OpTypeInt 16
-; CHECK-SPIRV: %[[#short2:]] = OpTypeVector %[[#short]] 2
-; CHECK-SPIRV: OpBitReverse %[[#short2]]
-
-define spir_kernel void @testBitRev(<2 x i16> %a, <2 x i16> %b, <2 x i16> %c, <2 x i16> addrspace(1)* nocapture %res) local_unnamed_addr {
-entry:
- %call = tail call <2 x i16> @llvm.bitreverse.v2i16(<2 x i16> %b)
- store <2 x i16> %call, <2 x i16> addrspace(1)* %res, align 4
- ret void
-}
-
-declare <2 x i16> @llvm.bitreverse.v2i16(<2 x i16>)
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_v2i32.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_v2i32.ll
new file mode 100644
index 0000000000000..3e2ed8b373e4d
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_v2i32.ll
@@ -0,0 +1,20 @@
+; RUN: llc -O0 -mtriple=spirv-unknown-vulkan %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - -filetype=obj | spirv-val --target-env vulkan1.3 %}
+
+; CHECK-SPIRV: %[[#short:]] = OpTypeInt 32
+; CHECK-SPIRV: %[[#short2:]] = OpTypeVector %[[#short]] 2
+; CHECK-SPIRV: OpBitReverse %[[#short2]]
+
+define hidden spir_func void @testBitRev(<2 x i32> %a, <2 x i32> %b, <2 x i32> %c, ptr %res) {
+entry:
+ %call = tail call <2 x i32> @llvm.bitreverse.v2i32(<2 x i32> %b)
+ store <2 x i32> %call, ptr %res, align 4
+ ret void
+}
+
+define void @main() #1 {
+ ret void
+}
+
+declare <2 x i32> @llvm.bitreverse.v2i32(<2 x i32>)
+attributes #1 = { "hlsl.numthreads"="8,1,1" "hlsl.shader"="compute" }
diff --git a/llvm/test/CodeGen/SPIRV/zero-length-array.ll b/llvm/test/CodeGen/SPIRV/zero-length-array.ll
index 668bf2018dec7..666176c87adb6 100644
--- a/llvm/test/CodeGen/SPIRV/zero-length-array.ll
+++ b/llvm/test/CodeGen/SPIRV/zero-length-array.ll
@@ -1,9 +1,8 @@
-; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
-; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - -filetype=obj | spirv-val %}
; CHECK: %[[#type:]] = OpTypeInt 32 0
-; CHECK: %[[#ext:]] = OpTypeRuntimeArray %[[#type]]
-; CHECK: %[[#]] = OpTypePointer Function %[[#ext]]
+; CHECK: %[[#ext:]] = OpConstant %[[#type]] 0
define spir_func void @_Z3foov() {
entry:
More information about the llvm-commits
mailing list