Description
The following slang shader does compile to CUDA without error, although it uses unimplemented texture subscript operator.
Not sure if this is an isolated occurence, but Slang really should report an error if the code uses unsupported features.
[shader("compute")]
[numthreads(16, 16, 1)]
void count(uint3 tid: SV_DispatchThreadID, Texture2D<float4> render_texture, RWByteAddressBuffer count_buffer)
{
if (any(tid.xy >= uint2(16)))
return;
float4 val = render_texture[tid.xy];
count_buffer.InterlockedAdd(0, val.x > 0 ? 1 : 0);
count_buffer.InterlockedAdd(4, val.y > 0 ? 1 : 0);
count_buffer.InterlockedAdd(8, val.z > 0 ? 1 : 0);
count_buffer.InterlockedAdd(12, val.w > 0 ? 1 : 0);
}The generate CUDA code is:
#include "C:/src/slang-rhi/build/_deps/slang-src/include/slang-cuda-prelude.h"
#line 6545 "hlsl.meta.slang"
__device__ bool any_0(bool2 x_0)
{
#line 6545
bool result_0 = false;
#line 6545
int i_0 = int(0);
#line 6585
for(;;)
{
#line 6585
if(i_0 < int(2))
{
}
else
{
#line 6585
break;
}
#line 6586
if(result_0)
{
#line 6586
result_0 = true;
#line 6586
}
else
{
#line 6586
result_0 = (bool((_slang_vector_get_element(x_0, i_0))));
#line 6586
}
#line 6585
i_0 = i_0 + int(1);
#line 6585
}
return result_0;
}
#line 3 "test.slang"
extern "C" __global__ void count(CUtexObject render_texture_0, RWByteAddressBuffer count_buffer_0)
{
int2 _S1 = make_int2 (int(16));
#line 5
uint2 _S2 = make_uint2 ((uint)_S1.x, (uint)_S1.y);
#line 5
if(any_0(uint2 {(blockIdx * blockDim + threadIdx).x, (blockIdx * blockDim + threadIdx).y} >= _S2))
{
#line 6
return;
}
#line 7
}