add the HIPRT_BAKE_KERNEL_GENERATED flag

Author: RichardGe

CMakeLists.txt

@@ -2,14 +2,26 @@ cmake_minimum_required(VERSION 3.10)
 project(hiprt)
 
 # Options
-option(BAKE_KERNEL "Encrypt and bake kernels" OFF)
+option(BAKE_KERNEL "enable the use of encrypted and baked kernels" OFF)
 option(BITCODE "Enable bitcode linking" OFF)
 option(PRECOMPILE "Precompile kernels" OFF)
 option(HIPRTEW "Use hiprtew" OFF)
 option(NO_ENCRYPT "Don't encrypt kernel source and binaries" OFF)
 option(NO_UNITTEST "Don't build unit tests" OFF)
 option(HIPRT_PREFER_HIP_5 "Prefer HIP 5" OFF)
 
+
+# GENERATE_BAKE_KERNEL is enabled by default if we use the flags 'BAKE_KERNEL' or 'BITCODE'.
+# It can be forced to OFF, but in this case, some building functions from the HIPRT API, like hiprtBuildTraceKernelsFromBitcode will fail.
+if(BAKE_KERNEL OR BITCODE)
+	set(GENERATE_BAKE_KERNEL___DEFAULT ON)
+else()
+	set(GENERATE_BAKE_KERNEL___DEFAULT OFF)
+endif()
+option(GENERATE_BAKE_KERNEL "generate the baked kernels" ${GENERATE_BAKE_KERNEL___DEFAULT})
+message(STATUS "GENERATE_BAKE_KERNEL= ${GENERATE_BAKE_KERNEL}")
+
+
 # Set C++ Standard
 set(CMAKE_CXX_STANDARD 17)
 set(CMAKE_CXX_STANDARD_REQUIRED ON)
@@ -322,7 +334,7 @@ if(PRECOMPILE)
 endif()
 
 
-if(BAKE_KERNEL OR BITCODE)
+if(BAKE_KERNEL OR GENERATE_BAKE_KERNEL)
 	message(">> BakeKernel Executed")
 	if(WIN32)
 		execute_process(COMMAND mkdir -p ${CMAKE_SOURCE_DIR}/hiprt/cache)
@@ -335,6 +347,9 @@ if(BAKE_KERNEL OR BITCODE)
 	if(BAKE_KERNEL)
 		target_compile_definitions(${HIPRT_NAME} PRIVATE HIPRT_LOAD_FROM_STRING ORO_PP_LOAD_FROM_STRING)
 	endif()
+	
+	target_compile_definitions(${HIPRT_NAME} PRIVATE HIPRT_BAKE_KERNEL_GENERATED)
+	
 endif()
 
 if(WIN32)

hiprt/hiprt_common.h

@@ -93,7 +93,7 @@
 #define HIPRT_HOST_DEVICE __host__ __device__
 
 // TODO: cleanup after baking is removed
-#if defined( HIPRT_LOAD_FROM_STRING ) || defined( HIPRT_BITCODE_LINKING )
+#if defined( HIPRT_BAKE_KERNEL_GENERATED )
 #define GET_ARGS( X ) ( hip::X##Args )
 #define GET_INC( X ) ( hip::X##Includes )
 #else

hiprt/impl/Compiler.cpp

@@ -31,7 +31,7 @@
 #include <contrib/easy-encryption/encrypt.h>
 #endif
 #include <regex>
-#if defined( HIPRT_LOAD_FROM_STRING ) || defined( HIPRT_BITCODE_LINKING )
+#if defined( HIPRT_BAKE_KERNEL_GENERATED )
 #include <hiprt/cache/Kernels.h>
 #include <hiprt/cache/KernelArgs.h>
 #endif
@@ -49,6 +49,12 @@ constexpr auto UseBakedCode = true;
 #else
 constexpr auto UseBakedCode = false;
 #endif
+
+#if defined( HIPRT_BAKE_KERNEL_GENERATED )
+constexpr auto BakedCodeIsGenerated = true;
+#else
+constexpr auto BakedCodeIsGenerated = false;
+#endif
 } // namespace
 
 namespace hiprt
@@ -224,7 +230,7 @@ void Compiler::buildKernels(
 			std::string				 extSrc = src;
 			if ( extended )
 			{
-				if constexpr ( UseBakedCode || UseBitcode )
+				if constexpr ( BakedCodeIsGenerated )
 				{
 					extSrc = "#include <hiprt_device_impl.h>\n";
 					addCustomFuncsSwitchCase( extSrc, funcNameSets, numGeomTypes, numRayTypes );
@@ -776,7 +782,7 @@ oroFunction Compiler::getFunctionFromPrecompiledBinary( const std::string& funcN
 std::string Compiler::buildFunctionTableBitcode(
 	Context& context, uint32_t numGeomTypes, uint32_t numRayTypes, const std::vector<hiprtFuncNameSet>& funcNameSets )
 {
-	if constexpr ( UseBitcode )
+	if constexpr ( BakedCodeIsGenerated )
 	{
 		bool amd = oroGetCurAPI( 0 ) == ORO_API_HIP;
 
@@ -847,7 +853,7 @@ std::string Compiler::buildFunctionTableBitcode(
 	}
 	else
 	{
-		throw std::runtime_error( "Not compiled with the bitcode linking support" );
+		throw std::runtime_error( "Not compiled with the baked kernel code support" );
 	}
 }
 } // namespace hiprt

premake5.lua

@@ -251,6 +251,7 @@ workspace "hiprt"
             defines {"HIPRT_LOAD_FROM_STRING"}
             defines { "ORO_PP_LOAD_FROM_STRING" }
         end
+        defines {"HIPRT_BAKE_KERNEL_GENERATED"}
     end
     if os.istarget("windows") then
         links{ "version" }