From 4e2f2771eb2f8991014d957848a6a25aa49c0aaf Mon Sep 17 00:00:00 2001 From: Yong He Date: Thu, 4 Jun 2020 12:42:17 -0700 Subject: Emit [loop] attribute to output HLSL. --- tests/cross-compile/loop-attribs.slang | 19 ++++++++ tests/cross-compile/loop-attribs.slang.expected | 61 +++++++++++++++++++++++++ 2 files changed, 80 insertions(+) create mode 100644 tests/cross-compile/loop-attribs.slang create mode 100644 tests/cross-compile/loop-attribs.slang.expected (limited to 'tests') diff --git a/tests/cross-compile/loop-attribs.slang b/tests/cross-compile/loop-attribs.slang new file mode 100644 index 000000000..ddedb67df --- /dev/null +++ b/tests/cross-compile/loop-attribs.slang @@ -0,0 +1,19 @@ +// loop-attribs.slang +// Test that loop attributes are correctly emitted to the resulting HLSL. + +//TEST:SIMPLE:-target hlsl -entry main -stage fragment -profile sm_6_0 + +float4 main() : SV_Target +{ + float sum = 0.0f; + + [loop] + for (int i = 0; i < 100; i++) + sum += float(i); + + [unroll(10)] + for (int j = 0; j < 100; j++) + sum += float(j); + + return float4(sum, 0, 0, 0); +} \ No newline at end of file diff --git a/tests/cross-compile/loop-attribs.slang.expected b/tests/cross-compile/loop-attribs.slang.expected new file mode 100644 index 000000000..0ecf0194f --- /dev/null +++ b/tests/cross-compile/loop-attribs.slang.expected @@ -0,0 +1,61 @@ +result code = 0 +standard error = { +} +standard output = { +#pragma pack_matrix(column_major) + +#line 6 "tests/cross-compile/loop-attribs.slang" +vector main() : SV_TARGET +{ + int i_0; + float sum_0; + int j_0; + float sum_1; + i_0 = int(0); + sum_0 = 0.00000000000000000000; + [loop] + for(;;) + { + +#line 11 + if(i_0 < int(100)) + { + } + else + { + break; + } + float _S1 = sum_0 + (float) i_0; + +#line 11 + int _S2 = i_0 + (int) int(1); + i_0 = _S2; + sum_0 = _S1; + } + j_0 = int(0); + sum_1 = sum_0; + [unroll] + for(;;) + { + +#line 15 + if(j_0 < int(100)) + { + } + else + { + break; + } + float _S3 = sum_1 + (float) j_0; + +#line 15 + int _S4 = j_0 + (int) int(1); + j_0 = _S4; + sum_1 = _S3; + } + +#line 18 + return vector(sum_1, (float) int(0), (float) int(0), (float) int(0)); +} + +} -- cgit v1.2.3 From 5e52b339f3e021cf8b68be56d1f3ff1d548a038f Mon Sep 17 00:00:00 2001 From: Yong He Date: Thu, 4 Jun 2020 14:04:54 -0700 Subject: Change loop-attrib test case to CROSS_COMPILE per review comments --- tests/cross-compile/loop-attribs.slang | 2 +- tests/cross-compile/loop-attribs.slang.expected | 61 ------------------------- 2 files changed, 1 insertion(+), 62 deletions(-) delete mode 100644 tests/cross-compile/loop-attribs.slang.expected (limited to 'tests') diff --git a/tests/cross-compile/loop-attribs.slang b/tests/cross-compile/loop-attribs.slang index ddedb67df..4336698c0 100644 --- a/tests/cross-compile/loop-attribs.slang +++ b/tests/cross-compile/loop-attribs.slang @@ -1,7 +1,7 @@ // loop-attribs.slang // Test that loop attributes are correctly emitted to the resulting HLSL. -//TEST:SIMPLE:-target hlsl -entry main -stage fragment -profile sm_6_0 +//TEST:CROSS_COMPILE:-target dxil-assembly -entry main -stage fragment -profile sm_6_0 float4 main() : SV_Target { diff --git a/tests/cross-compile/loop-attribs.slang.expected b/tests/cross-compile/loop-attribs.slang.expected deleted file mode 100644 index 0ecf0194f..000000000 --- a/tests/cross-compile/loop-attribs.slang.expected +++ /dev/null @@ -1,61 +0,0 @@ -result code = 0 -standard error = { -} -standard output = { -#pragma pack_matrix(column_major) - -#line 6 "tests/cross-compile/loop-attribs.slang" -vector main() : SV_TARGET -{ - int i_0; - float sum_0; - int j_0; - float sum_1; - i_0 = int(0); - sum_0 = 0.00000000000000000000; - [loop] - for(;;) - { - -#line 11 - if(i_0 < int(100)) - { - } - else - { - break; - } - float _S1 = sum_0 + (float) i_0; - -#line 11 - int _S2 = i_0 + (int) int(1); - i_0 = _S2; - sum_0 = _S1; - } - j_0 = int(0); - sum_1 = sum_0; - [unroll] - for(;;) - { - -#line 15 - if(j_0 < int(100)) - { - } - else - { - break; - } - float _S3 = sum_1 + (float) j_0; - -#line 15 - int _S4 = j_0 + (int) int(1); - j_0 = _S4; - sum_1 = _S3; - } - -#line 18 - return vector(sum_1, (float) int(0), (float) int(0), (float) int(0)); -} - -} -- cgit v1.2.3 From 899824eb3dbb327f3fd8479ca9bc01f8ea49ed98 Mon Sep 17 00:00:00 2001 From: Yong He Date: Thu, 4 Jun 2020 14:07:30 -0700 Subject: Add missing loop-attribs.slang.hlsl for the test case --- tests/cross-compile/loop-attribs.slang.hlsl | 55 +++++++++++++++++++++++++++++ 1 file changed, 55 insertions(+) create mode 100644 tests/cross-compile/loop-attribs.slang.hlsl (limited to 'tests') diff --git a/tests/cross-compile/loop-attribs.slang.hlsl b/tests/cross-compile/loop-attribs.slang.hlsl new file mode 100644 index 000000000..5d53f51e0 --- /dev/null +++ b/tests/cross-compile/loop-attribs.slang.hlsl @@ -0,0 +1,55 @@ +#pragma pack_matrix(column_major) + +#line 6 "tests/cross-compile/loop-attribs.slang" +vector main() : SV_TARGET +{ + int i_0; + float sum_0; + int j_0; + float sum_1; + i_0 = int(0); + sum_0 = 0.00000000000000000000; + [loop] + for(;;) + { + +#line 11 + if(i_0 < int(100)) + { + } + else + { + break; + } + float _S1 = sum_0 + (float) i_0; + +#line 11 + int _S2 = i_0 + (int) int(1); + i_0 = _S2; + sum_0 = _S1; + } + j_0 = int(0); + sum_1 = sum_0; + [unroll] + for(;;) + { + +#line 15 + if(j_0 < int(100)) + { + } + else + { + break; + } + float _S3 = sum_1 + (float) j_0; + +#line 15 + int _S4 = j_0 + (int) int(1); + j_0 = _S4; + sum_1 = _S3; + } + +#line 18 + return vector(sum_1, (float) int(0), (float) int(0), (float) int(0)); +} \ No newline at end of file -- cgit v1.2.3 From 3bb780724830ae830657a47e4eba008a4c0f4ff7 Mon Sep 17 00:00:00 2001 From: Tim Foley Date: Fri, 5 Jun 2020 10:00:15 -0700 Subject: Fixes for active mask synthesis + tests (#1370) * Fixes for active mask synthesis + tests There are two fixes here: * The code generation that follows active mask synthesis was requiring CUDA SM architecture version 7.0 for one of the introduced instructions, but not all of them. This change centralizes the handling of upgrading the required CUDA SM architecture version, and makes sure that the instructions introduced by active mask synthesis request version 7.0. * The tests for active mask synthesis were not flagged as requiring the `cuda_sm_7_0` feature when invoking `render-test-tool`, which meant they run but produce unexpected results when invoked on a GPU without the required semantics for functions like `__ballot_sync()`. This change adds the missing `-render-feature cuda_sm_7_0` to those tests. * fixup: mark more tests that rely on implicit active mask --- source/slang/slang-emit-cuda.cpp | 23 +++++++++++----------- source/slang/slang-emit-cuda.h | 3 +++ tests/hlsl-intrinsic/active-mask/for-break.slang | 2 +- .../active-mask/for-continue-ext.slang | 2 +- .../hlsl-intrinsic/active-mask/for-continue.slang | 2 +- tests/hlsl-intrinsic/active-mask/for.slang | 2 +- .../active-mask/if-conditional-exit.slang | 2 +- .../hlsl-intrinsic/active-mask/if-early-exit.slang | 2 +- .../hlsl-intrinsic/active-mask/if-one-sided.slang | 2 +- tests/hlsl-intrinsic/active-mask/if.slang | 2 +- .../active-mask/switch-no-default.slang | 2 +- .../active-mask/switch-trivial-fallthrough.slang | 2 +- tests/hlsl-intrinsic/active-mask/switch.slang | 7 ++++++- tests/hlsl-intrinsic/wave-active-product.slang | 2 +- tests/hlsl-intrinsic/wave-broadcast-lane-at.slang | 2 +- tests/hlsl-intrinsic/wave-diverge.slang | 2 +- tests/hlsl-intrinsic/wave-is-first-lane.slang | 2 +- tests/hlsl-intrinsic/wave-matrix.slang | 2 +- tests/hlsl-intrinsic/wave-prefix-product.slang | 2 +- tests/hlsl-intrinsic/wave-prefix-sum.slang | 2 +- tests/hlsl-intrinsic/wave-read-lane-at.slang | 2 +- tests/hlsl-intrinsic/wave-shuffle-vk.slang | 2 +- tests/hlsl-intrinsic/wave-shuffle.slang | 2 +- tests/hlsl-intrinsic/wave-vector.slang | 2 +- tests/hlsl-intrinsic/wave.slang | 2 +- 25 files changed, 43 insertions(+), 34 deletions(-) (limited to 'tests') 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 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 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 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 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 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 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 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 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 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 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 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 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 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 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 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 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 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 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 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 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 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 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 outputBuffer; -- cgit v1.2.3