diff options
25 files changed, 43 insertions, 34 deletions
diff --git a/source/slang/slang-emit-cuda.cpp b/source/slang/slang-emit-cuda.cpp index 639e7f737..25b06027d 100644 --- a/source/slang/slang-emit-cuda.cpp +++ b/source/slang/slang-emit-cuda.cpp @@ -520,6 +520,8 @@ bool CUDASourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOu } case kIROp_WaveMaskBallot: { + _requireCUDASMVersion(SemanticVersion(7, 0)); + m_writer->emit("__ballot_sync("); emitOperand(inst->getOperand(0), getInfo(EmitOp::General)); m_writer->emit(", "); @@ -529,12 +531,7 @@ bool CUDASourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOu } case kIROp_WaveMaskMatch: { - SemanticVersion version; - version.set(7, 0); - if (version > m_extensionTracker->m_smVersion) - { - m_extensionTracker->m_smVersion = version; - } + _requireCUDASMVersion(SemanticVersion(7, 0)); m_writer->emit("__match_any_sync("); emitOperand(inst->getOperand(0), getInfo(EmitOp::General)); @@ -549,6 +546,14 @@ bool CUDASourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOu return Super::tryEmitInstExprImpl(inst, inOuterPrec); } +void CUDASourceEmitter::_requireCUDASMVersion(SemanticVersion const& version) +{ + if (version > m_extensionTracker->m_smVersion) + { + m_extensionTracker->m_smVersion = version; + } +} + void CUDASourceEmitter::handleCallExprDecorationsImpl(IRInst* funcValue) { // Does this function declare any requirements on GLSL version or @@ -566,11 +571,7 @@ void CUDASourceEmitter::handleCallExprDecorationsImpl(IRInst* funcValue) { SemanticVersion version; version.setFromInteger(SemanticVersion::IntegerType(smDecoration->getCUDASMVersion())); - - if (version > m_extensionTracker->m_smVersion) - { - m_extensionTracker->m_smVersion = version; - } + _requireCUDASMVersion(version); } } } diff --git a/source/slang/slang-emit-cuda.h b/source/slang/slang-emit-cuda.h index 669bf2d20..c0a5ac5dc 100644 --- a/source/slang/slang-emit-cuda.h +++ b/source/slang/slang-emit-cuda.h @@ -79,6 +79,9 @@ protected: void _emitInitializerList(IRType* elementType, IRUse* operands, Index operandCount); void _emitInitializerListValue(IRType* elementType, IRInst* value); + /// Ensure that the generated code is compiled for at least CUDA SM `version` + void _requireCUDASMVersion(SemanticVersion const& version); + RefPtr<CUDAExtensionTracker> m_extensionTracker; }; diff --git a/tests/hlsl-intrinsic/active-mask/for-break.slang b/tests/hlsl-intrinsic/active-mask/for-break.slang index d0c35a017..440e7481c 100644 --- a/tests/hlsl-intrinsic/active-mask/for-break.slang +++ b/tests/hlsl-intrinsic/active-mask/for-break.slang @@ -8,7 +8,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 -xslang -DHACK //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -xslang -DHACK -//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name buffer RWStructuredBuffer<int> buffer; diff --git a/tests/hlsl-intrinsic/active-mask/for-continue-ext.slang b/tests/hlsl-intrinsic/active-mask/for-continue-ext.slang index 67ee536eb..881c0389e 100644 --- a/tests/hlsl-intrinsic/active-mask/for-continue-ext.slang +++ b/tests/hlsl-intrinsic/active-mask/for-continue-ext.slang @@ -9,7 +9,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 -xslang -DHACK //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -xslang -DHACK -//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name buffer RWStructuredBuffer<int> buffer; diff --git a/tests/hlsl-intrinsic/active-mask/for-continue.slang b/tests/hlsl-intrinsic/active-mask/for-continue.slang index 6bb0f23fb..7700a937a 100644 --- a/tests/hlsl-intrinsic/active-mask/for-continue.slang +++ b/tests/hlsl-intrinsic/active-mask/for-continue.slang @@ -9,7 +9,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 -xslang -DHACK //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -xslang -DHACK -//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name buffer RWStructuredBuffer<int> buffer; diff --git a/tests/hlsl-intrinsic/active-mask/for.slang b/tests/hlsl-intrinsic/active-mask/for.slang index aac66f3c9..c6fef5e58 100644 --- a/tests/hlsl-intrinsic/active-mask/for.slang +++ b/tests/hlsl-intrinsic/active-mask/for.slang @@ -8,7 +8,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name buffer RWStructuredBuffer<int> buffer; diff --git a/tests/hlsl-intrinsic/active-mask/if-conditional-exit.slang b/tests/hlsl-intrinsic/active-mask/if-conditional-exit.slang index 78815e942..9f90b9380 100644 --- a/tests/hlsl-intrinsic/active-mask/if-conditional-exit.slang +++ b/tests/hlsl-intrinsic/active-mask/if-conditional-exit.slang @@ -9,7 +9,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name buffer diff --git a/tests/hlsl-intrinsic/active-mask/if-early-exit.slang b/tests/hlsl-intrinsic/active-mask/if-early-exit.slang index 8532035a3..35efb9a6d 100644 --- a/tests/hlsl-intrinsic/active-mask/if-early-exit.slang +++ b/tests/hlsl-intrinsic/active-mask/if-early-exit.slang @@ -8,7 +8,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name buffer RWStructuredBuffer<int> buffer; diff --git a/tests/hlsl-intrinsic/active-mask/if-one-sided.slang b/tests/hlsl-intrinsic/active-mask/if-one-sided.slang index 0662fc1e7..9b3259f3b 100644 --- a/tests/hlsl-intrinsic/active-mask/if-one-sided.slang +++ b/tests/hlsl-intrinsic/active-mask/if-one-sided.slang @@ -6,7 +6,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name buffer RWStructuredBuffer<int> buffer; diff --git a/tests/hlsl-intrinsic/active-mask/if.slang b/tests/hlsl-intrinsic/active-mask/if.slang index 897ba6d1b..077bb4b40 100644 --- a/tests/hlsl-intrinsic/active-mask/if.slang +++ b/tests/hlsl-intrinsic/active-mask/if.slang @@ -6,7 +6,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name buffer RWStructuredBuffer<int> buffer; diff --git a/tests/hlsl-intrinsic/active-mask/switch-no-default.slang b/tests/hlsl-intrinsic/active-mask/switch-no-default.slang index d103bc4c7..f3060631b 100644 --- a/tests/hlsl-intrinsic/active-mask/switch-no-default.slang +++ b/tests/hlsl-intrinsic/active-mask/switch-no-default.slang @@ -9,7 +9,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 -xslang -DHACK //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -xslang -DHACK -//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name buffer RWStructuredBuffer<int> buffer; diff --git a/tests/hlsl-intrinsic/active-mask/switch-trivial-fallthrough.slang b/tests/hlsl-intrinsic/active-mask/switch-trivial-fallthrough.slang index ceb7d236d..b5674c8d1 100644 --- a/tests/hlsl-intrinsic/active-mask/switch-trivial-fallthrough.slang +++ b/tests/hlsl-intrinsic/active-mask/switch-trivial-fallthrough.slang @@ -13,7 +13,7 @@ // target because we do not synthesize the active // mask value we want/expect to see. // -//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name buffer RWStructuredBuffer<int> buffer; diff --git a/tests/hlsl-intrinsic/active-mask/switch.slang b/tests/hlsl-intrinsic/active-mask/switch.slang index 41f379f3d..0ab0c7173 100644 --- a/tests/hlsl-intrinsic/active-mask/switch.slang +++ b/tests/hlsl-intrinsic/active-mask/switch.slang @@ -6,7 +6,12 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 -xslang -DHACK //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -xslang -DHACK -//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute + +// Note: this test is currently disabled on the CUDA +// target because we do not synthesize the active +// mask value we want/expect to see. +// +//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name buffer RWStructuredBuffer<int> buffer; diff --git a/tests/hlsl-intrinsic/wave-active-product.slang b/tests/hlsl-intrinsic/wave-active-product.slang index ca3fdcb77..67b9ca587 100644 --- a/tests/hlsl-intrinsic/wave-active-product.slang +++ b/tests/hlsl-intrinsic/wave-active-product.slang @@ -2,7 +2,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name outputBuffer RWStructuredBuffer<int> outputBuffer; diff --git a/tests/hlsl-intrinsic/wave-broadcast-lane-at.slang b/tests/hlsl-intrinsic/wave-broadcast-lane-at.slang index b6f5d3847..bb1523219 100644 --- a/tests/hlsl-intrinsic/wave-broadcast-lane-at.slang +++ b/tests/hlsl-intrinsic/wave-broadcast-lane-at.slang @@ -3,7 +3,7 @@ //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 // Disabled on VK because glsl can't do WaveReadLaneAt on matrix. //DISABLE_TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name outputBuffer RWStructuredBuffer<int> outputBuffer; diff --git a/tests/hlsl-intrinsic/wave-diverge.slang b/tests/hlsl-intrinsic/wave-diverge.slang index ab83a1553..2d500d0c8 100644 --- a/tests/hlsl-intrinsic/wave-diverge.slang +++ b/tests/hlsl-intrinsic/wave-diverge.slang @@ -2,7 +2,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name outputBuffer RWStructuredBuffer<int> outputBuffer; diff --git a/tests/hlsl-intrinsic/wave-is-first-lane.slang b/tests/hlsl-intrinsic/wave-is-first-lane.slang index 39e7dab5d..e1f0faef3 100644 --- a/tests/hlsl-intrinsic/wave-is-first-lane.slang +++ b/tests/hlsl-intrinsic/wave-is-first-lane.slang @@ -2,7 +2,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//TEST(compute, vulkan):COMPARE_COMPUTE_EX:-cuda -compute +//TEST(compute, vulkan):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name outputBuffer RWStructuredBuffer<int> outputBuffer; diff --git a/tests/hlsl-intrinsic/wave-matrix.slang b/tests/hlsl-intrinsic/wave-matrix.slang index b5af69f5d..50acc4a71 100644 --- a/tests/hlsl-intrinsic/wave-matrix.slang +++ b/tests/hlsl-intrinsic/wave-matrix.slang @@ -2,7 +2,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //DISABLE_TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0], stride=4):out,name outputBuffer RWStructuredBuffer<int> outputBuffer; diff --git a/tests/hlsl-intrinsic/wave-prefix-product.slang b/tests/hlsl-intrinsic/wave-prefix-product.slang index a56912616..436ce06f7 100644 --- a/tests/hlsl-intrinsic/wave-prefix-product.slang +++ b/tests/hlsl-intrinsic/wave-prefix-product.slang @@ -2,7 +2,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0], stride=4):out,name outputBuffer RWStructuredBuffer<int> outputBuffer; diff --git a/tests/hlsl-intrinsic/wave-prefix-sum.slang b/tests/hlsl-intrinsic/wave-prefix-sum.slang index 343a7afbd..358debedc 100644 --- a/tests/hlsl-intrinsic/wave-prefix-sum.slang +++ b/tests/hlsl-intrinsic/wave-prefix-sum.slang @@ -2,7 +2,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0], stride=4):out,name outputBuffer RWStructuredBuffer<int> outputBuffer; diff --git a/tests/hlsl-intrinsic/wave-read-lane-at.slang b/tests/hlsl-intrinsic/wave-read-lane-at.slang index c3caaa4e8..4d840ad0f 100644 --- a/tests/hlsl-intrinsic/wave-read-lane-at.slang +++ b/tests/hlsl-intrinsic/wave-read-lane-at.slang @@ -3,7 +3,7 @@ //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 // Disabled on VK because glsl can't do WaveReadLaneAt on matrix. //DISABLE_TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name outputBuffer RWStructuredBuffer<int> outputBuffer; diff --git a/tests/hlsl-intrinsic/wave-shuffle-vk.slang b/tests/hlsl-intrinsic/wave-shuffle-vk.slang index 75aa392ea..820a69e89 100644 --- a/tests/hlsl-intrinsic/wave-shuffle-vk.slang +++ b/tests/hlsl-intrinsic/wave-shuffle-vk.slang @@ -3,7 +3,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name outputBuffer RWStructuredBuffer<int> outputBuffer; diff --git a/tests/hlsl-intrinsic/wave-shuffle.slang b/tests/hlsl-intrinsic/wave-shuffle.slang index 093babcce..608e57ffd 100644 --- a/tests/hlsl-intrinsic/wave-shuffle.slang +++ b/tests/hlsl-intrinsic/wave-shuffle.slang @@ -4,7 +4,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 // Disabled because vk doesn't currently support matrix types. See wave-shuffle-vk.slang //DISABLE_TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name outputBuffer RWStructuredBuffer<int> outputBuffer; diff --git a/tests/hlsl-intrinsic/wave-vector.slang b/tests/hlsl-intrinsic/wave-vector.slang index 8d2868600..784f16c16 100644 --- a/tests/hlsl-intrinsic/wave-vector.slang +++ b/tests/hlsl-intrinsic/wave-vector.slang @@ -2,7 +2,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0], stride=4):out,name outputBuffer RWStructuredBuffer<int> outputBuffer; diff --git a/tests/hlsl-intrinsic/wave.slang b/tests/hlsl-intrinsic/wave.slang index d8273080c..eb647d0fe 100644 --- a/tests/hlsl-intrinsic/wave.slang +++ b/tests/hlsl-intrinsic/wave.slang @@ -2,7 +2,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name outputBuffer RWStructuredBuffer<int> outputBuffer; |