diff options
Diffstat (limited to 'source/slang')
| -rw-r--r-- | source/slang/hlsl.meta.slang | 1 | ||||
| -rw-r--r-- | source/slang/slang-compiler.cpp | 50 | ||||
| -rw-r--r-- | source/slang/slang-emit-cuda.cpp | 60 | ||||
| -rw-r--r-- | source/slang/slang.vcxproj | 13 | ||||
| -rw-r--r-- | source/slang/slang.vcxproj.filters | 45 |
5 files changed, 126 insertions, 43 deletions
diff --git a/source/slang/hlsl.meta.slang b/source/slang/hlsl.meta.slang index 483c2299f..4279e4a4e 100644 --- a/source/slang/hlsl.meta.slang +++ b/source/slang/hlsl.meta.slang @@ -3162,6 +3162,7 @@ void AcceptHitAndEndSearch(); // 10.4.1 - Ray Dispatch System Values __target_intrinsic(glsl, "(gl_LaunchIDNV)") +__target_intrinsic(cuda, "optixGetLaunchIndex") uint3 DispatchRaysIndex(); __target_intrinsic(glsl, "(gl_LaunchSizeNV)") diff --git a/source/slang/slang-compiler.cpp b/source/slang/slang-compiler.cpp index 53a028483..f81a63275 100644 --- a/source/slang/slang-compiler.cpp +++ b/source/slang/slang-compiler.cpp @@ -1373,6 +1373,56 @@ SlangResult dissassembleDXILUsingDXC( default: SLANG_ASSERT(!"Unhandled floating point mode"); } + { + // We need to look at the stage of the entry point(s) we are + // being asked to compile, since this will determine the + // "pipeline" that the result should be compiled for (e.g., + // compute vs. ray tracing). + // + // TODO: This logic is kind of messy in that it assumes + // a program to be compiled will only contain kernels for + // a single pipeline type, but that invariant isn't expressed + // at all in the front-end today. It also has no error + // checking for the case where there are conflicts. + // + // HACK: Right now none of the above concerns matter + // because we always perform code generation on a single + // entry point at a time. + // + Index entryPointCount = slangRequest->getProgram()->getEntryPointCount(); + for(Index ee = 0; ee < entryPointCount; ++ee) + { + auto stage = slangRequest->getProgram()->getEntryPoint(ee)->getStage(); + switch(stage) + { + default: + break; + + case Stage::Compute: + options.pipelineType = DownstreamCompiler::PipelineType::Compute; + break; + + case Stage::Vertex: + case Stage::Hull: + case Stage::Domain: + case Stage::Geometry: + case Stage::Fragment: + options.pipelineType = DownstreamCompiler::PipelineType::Rasterization; + break; + + case Stage::RayGeneration: + case Stage::Intersection: + case Stage::AnyHit: + case Stage::ClosestHit: + case Stage::Miss: + case Stage::Callable: + options.pipelineType = DownstreamCompiler::PipelineType::RayTracing; + break; + } + } + + } + // Add all the search paths (as calculated earlier - they will only be set if this is a pass through else will be empty) options.includePaths = includePaths; diff --git a/source/slang/slang-emit-cuda.cpp b/source/slang/slang-emit-cuda.cpp index a728df755..702543fc8 100644 --- a/source/slang/slang-emit-cuda.cpp +++ b/source/slang/slang-emit-cuda.cpp @@ -606,21 +606,57 @@ void CUDASourceEmitter::emitModuleImpl(IRModule* module) IREntryPointDecoration* entryPointDecor = func->findDecoration<IREntryPointDecoration>(); - if (entryPointDecor && entryPointDecor->getProfile().GetStage() == Stage::Compute) + if (entryPointDecor) { - Int sizeAlongAxis[kThreadGroupAxisCount]; - getComputeThreadGroupSize(func, sizeAlongAxis); - - // - m_writer->emit("// [numthreads("); - for (int ii = 0; ii < kThreadGroupAxisCount; ++ii) + // We have an entry-point function in the IR module, which we + // will want to emit as a `__global__` function in the generated + // CUDA C++. + // + // The most common case will be a compute kernel, in which case + // we will emit the function more or less as-is, including + // usingits original name as the name of the global symbol. + // + String funcName = getName(func); + String globalSymbolName = funcName; + + // We also suport emitting ray tracing kernels for use with + // OptiX, and in that case the name of the global symbol + // must be prefixed to indicate to the OptiX runtime what + // stage it is to be compiled for. + // + auto stage = entryPointDecor->getProfile().GetStage(); + switch( stage ) { - if (ii != 0) m_writer->emit(", "); - m_writer->emit(sizeAlongAxis[ii]); + default: + break; + + #define CASE(STAGE, PREFIX) \ + case Stage::STAGE: globalSymbolName = #PREFIX + funcName; break + + CASE(RayGeneration, __raygen__); + // TODO: Add the other ray tracing shader stages here. + #undef CASE } - m_writer->emit(")]\n"); - String funcName = getName(func); + // As a convenience for anybody reading the generated + // CUDA C++ code, we will prefix a compute kernel + // with the information from the `[numthreads(...)]` + // attribute in the source. + // + if(stage == Stage::Compute) + { + Int sizeAlongAxis[kThreadGroupAxisCount]; + getComputeThreadGroupSize(func, sizeAlongAxis); + + // + m_writer->emit("// [numthreads("); + for (int ii = 0; ii < kThreadGroupAxisCount; ++ii) + { + if (ii != 0) m_writer->emit(", "); + m_writer->emit(sizeAlongAxis[ii]); + } + m_writer->emit(")]\n"); + } m_writer->emit("extern \"C\" __global__ "); @@ -628,7 +664,7 @@ void CUDASourceEmitter::emitModuleImpl(IRModule* module) // Emit the actual function emitEntryPointAttributes(func, entryPointDecor); - emitType(resultType, funcName); + emitType(resultType, globalSymbolName); m_writer->emit("(UniformEntryPointParams* params, UniformState* uniformState)"); emitSemantics(func); diff --git a/source/slang/slang.vcxproj b/source/slang/slang.vcxproj index 027c43a95..766893da3 100644 --- a/source/slang/slang.vcxproj +++ b/source/slang/slang.vcxproj @@ -188,7 +188,6 @@ <ItemGroup> <ClInclude Include="..\..\slang.h" /> <ClInclude Include="core.meta.slang.h" /> - <ClInclude Include="glsl.meta.slang.h" /> <ClInclude Include="hlsl.meta.slang.h" /> <ClInclude Include="slang-check-impl.h" /> <ClInclude Include="slang-check.h" /> @@ -345,17 +344,13 @@ <ClCompile Include="slang.cpp" /> </ItemGroup> <ItemGroup> - <None Include="..\core\core.natvis" /> - <None Include="slang.natvis" /> - </ItemGroup> - <ItemGroup> <CustomBuild Include="core.meta.slang"> <FileType>Document</FileType> <Command Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">"../../bin/windows-x86/debug/slang-generate" %(Identity)</Command> <Command Condition="'$(Configuration)|$(Platform)'=='Debug|x64'">"../../bin/windows-x64/debug/slang-generate" %(Identity)</Command> <Command Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">"../../bin/windows-x86/release/slang-generate" %(Identity)</Command> <Command Condition="'$(Configuration)|$(Platform)'=='Release|x64'">"../../bin/windows-x64/release/slang-generate" %(Identity)</Command> - <Outputs>%(Identity).h</Outputs> + <Outputs>../../core.meta.slang.h</Outputs> <Message>slang-generate %(Identity)</Message> <AdditionalInputs Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">../../bin/windows-x86/debug/slang-generate.exe</AdditionalInputs> <AdditionalInputs Condition="'$(Configuration)|$(Platform)'=='Debug|x64'">../../bin/windows-x64/debug/slang-generate.exe</AdditionalInputs> @@ -368,7 +363,7 @@ <Command Condition="'$(Configuration)|$(Platform)'=='Debug|x64'">"../../bin/windows-x64/debug/slang-generate" %(Identity)</Command> <Command Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">"../../bin/windows-x86/release/slang-generate" %(Identity)</Command> <Command Condition="'$(Configuration)|$(Platform)'=='Release|x64'">"../../bin/windows-x64/release/slang-generate" %(Identity)</Command> - <Outputs>%(Identity).h</Outputs> + <Outputs>../../hlsl.meta.slang.h</Outputs> <Message>slang-generate %(Identity)</Message> <AdditionalInputs Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">../../bin/windows-x86/debug/slang-generate.exe</AdditionalInputs> <AdditionalInputs Condition="'$(Configuration)|$(Platform)'=='Debug|x64'">../../bin/windows-x64/debug/slang-generate.exe</AdditionalInputs> @@ -377,6 +372,10 @@ </CustomBuild> </ItemGroup> <ItemGroup> + <Natvis Include="..\core\core.natvis" /> + <Natvis Include="slang.natvis" /> + </ItemGroup> + <ItemGroup> <ProjectReference Include="..\core\core.vcxproj"> <Project>{F9BE7957-8399-899E-0C49-E714FDDD4B65}</Project> </ProjectReference> diff --git a/source/slang/slang.vcxproj.filters b/source/slang/slang.vcxproj.filters index 9be567654..442f545c6 100644 --- a/source/slang/slang.vcxproj.filters +++ b/source/slang/slang.vcxproj.filters @@ -1,4 +1,4 @@ -<?xml version="1.0" encoding="utf-8"?> +<?xml version="1.0" encoding="utf-8"?> <Project ToolsVersion="4.0" xmlns="http://schemas.microsoft.com/developer/msbuild/2003"> <ItemGroup> <Filter Include="Header Files"> @@ -15,9 +15,6 @@ <ClInclude Include="core.meta.slang.h"> <Filter>Header Files</Filter> </ClInclude> - <ClInclude Include="glsl.meta.slang.h"> - <Filter>Header Files</Filter> - </ClInclude> <ClInclude Include="hlsl.meta.slang.h"> <Filter>Header Files</Filter> </ClInclude> @@ -99,6 +96,9 @@ <ClInclude Include="slang-ir-glsl-legalize.h"> <Filter>Header Files</Filter> </ClInclude> + <ClInclude Include="slang-ir-inline.h"> + <Filter>Header Files</Filter> + </ClInclude> <ClInclude Include="slang-ir-inst-defs.h"> <Filter>Header Files</Filter> </ClInclude> @@ -138,6 +138,9 @@ <ClInclude Include="slang-ir-string-hash.h"> <Filter>Header Files</Filter> </ClInclude> + <ClInclude Include="slang-ir-strip-witness-tables.h"> + <Filter>Header Files</Filter> + </ClInclude> <ClInclude Include="slang-ir-strip.h"> <Filter>Header Files</Filter> </ClInclude> @@ -246,12 +249,6 @@ <ClInclude Include="slang-visitor.h"> <Filter>Header Files</Filter> </ClInclude> - <ClInclude Include="slang-ir-strip-witness-tables.h"> - <Filter>Header Files</Filter> - </ClInclude> - <ClInclude Include="slang-ir-inline.h"> - <Filter>Header Files</Filter> - </ClInclude> </ItemGroup> <ItemGroup> <ClCompile Include="slang-check-conformance.cpp"> @@ -353,6 +350,9 @@ <ClCompile Include="slang-ir-glsl-legalize.cpp"> <Filter>Source Files</Filter> </ClCompile> + <ClCompile Include="slang-ir-inline.cpp"> + <Filter>Source Files</Filter> + </ClCompile> <ClCompile Include="slang-ir-legalize-types.cpp"> <Filter>Source Files</Filter> </ClCompile> @@ -389,6 +389,9 @@ <ClCompile Include="slang-ir-string-hash.cpp"> <Filter>Source Files</Filter> </ClCompile> + <ClCompile Include="slang-ir-strip-witness-tables.cpp"> + <Filter>Source Files</Filter> + </ClCompile> <ClCompile Include="slang-ir-strip.cpp"> <Filter>Source Files</Filter> </ClCompile> @@ -470,20 +473,6 @@ <ClCompile Include="slang.cpp"> <Filter>Source Files</Filter> </ClCompile> - <ClCompile Include="slang-ir-strip-witness-tables.cpp"> - <Filter>Source Files</Filter> - </ClCompile> - <ClCompile Include="slang-ir-inline.cpp"> - <Filter>Source Files</Filter> - </ClCompile> - </ItemGroup> - <ItemGroup> - <None Include="..\core\core.natvis"> - <Filter>Source Files</Filter> - </None> - <None Include="slang.natvis"> - <Filter>Source Files</Filter> - </None> </ItemGroup> <ItemGroup> <CustomBuild Include="core.meta.slang"> @@ -493,4 +482,12 @@ <Filter>Source Files</Filter> </CustomBuild> </ItemGroup> + <ItemGroup> + <Natvis Include="..\core\core.natvis"> + <Filter>Source Files</Filter> + </Natvis> + <Natvis Include="slang.natvis"> + <Filter>Source Files</Filter> + </Natvis> + </ItemGroup> </Project>
\ No newline at end of file |