[ORO-0] Some RTC apis.

Note: Google Test filter = *kernelExec*:*init*:*deviceprops*:*malloc* [==========] Running 6 tests from 1 test case. [----------] Global test environment set-up. [----------] 6 tests from OroTestBase [ RUN ] OroTestBase.init [ OK ] OroTestBase.init (91 ms) [ RUN ] OroTestBase.deviceprops executing on Intel(R) Arc(TM) A770 Graphics () 512 multiProcessors [ OK ] OroTestBase.deviceprops (6 ms) [ RUN ] OroTestBase.malloc [ OK ] OroTestBase.malloc (28 ms) [ RUN ] OroTestBase.kernelExec [ OK ] OroTestBase.kernelExec (81 ms) [ RUN ] OroTestBase.kernelExecPreCompiled 0: 123 1: 123 2: 123 3: 123 4: 123 5: 123 6: 123 7: 123 8: 123 9: 123 10: 123 11: 123 12: 123 13: 123 14: 123 15: 123 16: 123 17: 123 18: 123 19: 123 20: 123 21: 123 22: 123 23: 123 24: 123 25: 123 26: 123 27: 123 28: 123 29: 123 30: 123 31: 123 32: 123 33: 123 34: 123 35: 123 36: 123 37: 123 38: 123 39: 123 40: 123 41: 123 42: 123 43: 123 44: 123 45: 123 46: 123 47: 123 48: 123 49: 123 50: 123 51: 123 52: 123 53: 123 54: 123 55: 123 56: 123 57: 123 58: 123 59: 123 60: 123 61: 123 62: 123 63: 123 [ OK ] OroTestBase.kernelExecPreCompiled (126 ms) [ RUN ] OroTestBase.kernelExecPreCompiled1 executing on Intel(R) Arc(TM) A770 Graphics () 512 multiProcessors [ OK ] OroTestBase.kernelExecPreCompiled1 (93 ms) [----------] 6 tests from OroTestBase (427 ms total) [----------] Global test environment tear-down [==========] 6 tests from 1 test case ran. (428 ms total) [ PASSED ] 6 tests.
amdadvtech · Mar 7, 2023 · 3eaabdc · 3eaabdc
1 parent b148fbc
commit 3eaabdc
Show file tree

Hide file tree

Showing 4 changed files with 72 additions and 5 deletions.
diff --git a/Orochi/Orochi.cpp b/Orochi/Orochi.cpp
@@ -30,7 +30,12 @@
 
 #if defined( ORO_INTEL )
 #include <ze_api.h>
+#include <fstream>
 //https://spec.oneapi.io/level-zero/latest/api.html
+
+const char* getTempFilePath() { return "__x.cl"; }
+const char* getCompiledFilePath() { return "__x_Gen12LPdg1.spv"; }
+std::string getCompiledPathPostfix() { return "_Gen12LPdg1.spv"; }
 #endif
 
 std::unordered_map<void*, oroCtx> s_oroCtxs;
@@ -45,6 +50,7 @@ struct IntelContext
 	ze_context_handle_t m_ctxt = 0;
 	ze_device_handle_t m_device = 0;
 	ze_command_list_handle_t m_cmdList = 0;
+	ze_command_queue_handle_t m_queue = 0;
 };
 
 static IntelContext s_iCtxt;
@@ -583,7 +589,6 @@ oroError OROAPI oroCtxDestroy(oroCtx ctx)
 #if defined( ORO_INTEL )
 	if( s_api == ORO_API_INTEL )
 	{
-		TODO( "oroCtxDestroy" );
 		auto e = zeContextDestroy( s_iCtxt.m_ctxt );
 		if( e != ZE_RESULT_SUCCESS ) return oroErrorUnknown;
 
@@ -807,6 +812,8 @@ oroError OROAPI oroMemcpyDtoH(void* dstHost, oroDeviceptr srcDevice, size_t Byte
 	{
 		auto e = zeCommandListAppendMemoryCopy( s_iCtxt.m_cmdList, dstHost, (const void*)srcDevice, ByteCount, 0, 0, 0 );
 		if( e != ZE_RESULT_SUCCESS ) return (oroError)e;
+		e = (ze_result_t)oroStreamSynchronize( (oroStream)s_iCtxt.m_queue );
+		if( e != ZE_RESULT_SUCCESS ) return (oroError)e;
 		return oroSuccess;
 	}
 #endif
@@ -844,6 +851,8 @@ oroError OROAPI oroMemset(oroDeviceptr dstDevice, unsigned int ui, size_t N)
 	{
 		auto e = zeCommandListAppendMemoryFill( s_iCtxt.m_cmdList, (void*)dstDevice, &ui, sizeof( int ), N, 0, 0, 0 );
 		if( e != ZE_RESULT_SUCCESS ) return (oroError)e;
+		e = (ze_result_t)oroStreamSynchronize( (oroStream)s_iCtxt.m_queue );
+		if( e != ZE_RESULT_SUCCESS ) return (oroError)e;
 		return oroSuccess;
 	}
 #endif
@@ -974,18 +983,43 @@ orortcResult OROAPI orortcAddNameExpression( orortcProgram prog, const char* nam
 }
 orortcResult OROAPI orortcCompileProgram(orortcProgram prog, int numOptions, const char** options)
 {
+#if defined( ORO_INTEL )
+	if( s_api == ORO_API_INTEL )
+	{
+		std::string cmd = "ocloc -file " + std::string( getTempFilePath() );
+		cmd += " -device dg1";// -output " + std::string( getCompiledFilePath() );
+		system( cmd.c_str() );
+		return ORORTC_SUCCESS;
+	}
+#endif
 	__ORORTC_FUNC1( CompileProgram( (nvrtcProgram)prog, numOptions, options ),
 		CompileProgram( (hiprtcProgram)prog, numOptions, options ) );
 	return ORORTC_ERROR_INTERNAL_ERROR;
 }
 orortcResult OROAPI orortcCreateProgram(orortcProgram* prog, const char* src, const char* name, int numHeaders, const char** headers, const char** includeNames)
 {
+#if defined( ORO_INTEL )
+	if( s_api == ORO_API_INTEL )
+	{
+		std::ofstream file;
+		file.open( getTempFilePath() );
+		file << src;
+		file.close();
+		return ORORTC_SUCCESS;
+	}
+#endif
 	__ORORTC_FUNC1( CreateProgram( (nvrtcProgram*)prog, src, name, numHeaders, headers, includeNames ), 
 		CreateProgram( (hiprtcProgram*)prog, src, name, numHeaders, headers, includeNames ) );
 	return ORORTC_ERROR_INTERNAL_ERROR;
 }
 orortcResult OROAPI orortcDestroyProgram(orortcProgram* prog)
 {
+#if defined( ORO_INTEL )
+	if( s_api == ORO_API_INTEL )
+	{
+		return ORORTC_SUCCESS;
+	}
+#endif
 	__ORORTC_FUNC1( DestroyProgram( (nvrtcProgram*)prog), 
 		DestroyProgram( (hiprtcProgram*)prog ) );
 	return ORORTC_ERROR_INTERNAL_ERROR;
@@ -1021,12 +1055,38 @@ orortcResult OROAPI orortcGetBitcodeSize(orortcProgram prog, size_t* bitcodeSize
 }
 orortcResult OROAPI orortcGetCode(orortcProgram prog, char* code)
 {
+#if defined( ORO_INTEL )
+	if( s_api == ORO_API_INTEL )
+	{
+		std::ifstream f( getCompiledFilePath() );
+		const auto b = f.tellg();
+		f.seekg( 0, std::ios::end );
+		const auto e = f.tellg();
+		const auto s = e - b;
+		f.seekg( 0, std::ios::beg );
+		f.read( code, s );
+		f.close();
+		return ORORTC_SUCCESS;
+	}
+#endif
 	__ORORTC_FUNC1( GetPTX( (nvrtcProgram)prog, code ), 
 		GetCode( (hiprtcProgram)prog, code ) );
 	return ORORTC_ERROR_INTERNAL_ERROR;
 }
 orortcResult OROAPI orortcGetCodeSize(orortcProgram prog, size_t* codeSizeRet)
 {
+#if defined( ORO_INTEL )
+	if( s_api == ORO_API_INTEL )
+	{
+		std::ifstream f( getCompiledFilePath() );
+		const auto b = f.tellg();
+		f.seekg( 0, std::ios::end );
+		const auto e = f.tellg();
+		*codeSizeRet = e - b;
+		f.close();
+		return ORORTC_SUCCESS;
+	}
+#endif
 	__ORORTC_FUNC1( GetPTXSize( (nvrtcProgram)prog, codeSizeRet ), 
 		GetCodeSize( (hiprtcProgram)prog, codeSizeRet ) );
 	return ORORTC_ERROR_INTERNAL_ERROR;
@@ -1119,6 +1179,7 @@ oroError OROAPI oroStreamCreate( oroStream* stream )
 		cmdQueueDesc.mode = ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS;
 		e = zeCommandQueueCreate( s_iCtxt.m_ctxt, dev, &cmdQueueDesc, (ze_command_queue_handle_t*)stream );
 		if( e != ZE_RESULT_SUCCESS ) return (oroError)e;
+		s_iCtxt.m_queue = *(ze_command_queue_handle_t*)stream;
 
 		ze_command_list_handle_t command_list;
 		ze_command_list_desc_t cmdListDesc = {};
@@ -1140,7 +1201,7 @@ oroError OROAPI oroStreamSynchronize( oroStream hStream )
 #if defined( ORO_INTEL )
 	if( s_api == ORO_API_INTEL )
 	{
-		if( hStream == 0 ) return oroErrorUnknown;
+		if( hStream == 0 ) hStream = (oroStream)s_iCtxt.m_queue;
 		ze_command_queue_handle_t queue = (ze_command_queue_handle_t)hStream;
 		auto e = zeCommandListClose( s_iCtxt.m_cmdList );
 		if( e != ZE_RESULT_SUCCESS ) return (oroError)e;

diff --git a/Orochi/OrochiUtils.cpp b/Orochi/OrochiUtils.cpp
@@ -501,7 +501,7 @@ oroFunction OrochiUtils::getFunction( oroDevice device, const char* code, const
 		OrochiUtilsImpl::cacheBinaryToFile( codec, cacheFile );
 	}
 	oroModule module;
-	oroError ee = oroModuleLoadData( &module, codec.data() );
+	oroError ee = oroModuleLoadData( &module, codec.data(), codec.size() );
 	OROASSERT( ee == oroSuccess, 0 );
 	ee = oroModuleGetFunction( &function, module, funcName );
 	OROASSERT( ee == oroSuccess, 0 );

diff --git a/UnitTest/main.cpp b/UnitTest/main.cpp
@@ -96,7 +96,8 @@ TEST_F( OroTestBase, kernelExec )
 	int* a_device = nullptr;
 	OROCHECK( oroMalloc( (oroDeviceptr*)&a_device, sizeof( int ) ) );
 	OROCHECK( oroMemset( (oroDeviceptr)a_device, 0, sizeof( int ) ) );
-	oroFunction kernel = o.getFunctionFromFile( m_device, "../UnitTest/testKernel.h", "testKernel", 0 ); 
+	oroFunction kernel = o.getFunctionFromFile( m_device, 
+		( oroGetCurAPI( 0 ) != ORO_API_INTEL ) ? "../UnitTest/testKernel.h" : "../UnitTest/testKernel.cl", "testKernel", 0 ); 
 	const void* args[] = { &a_device };
 	OrochiUtils::launch1D( kernel, 64, args, 64 );
 	OrochiUtils::waitForCompletion();
@@ -134,7 +135,7 @@ TEST_F( OroTestBase, kernelExecPreCompiled )
 		loadFile( path.c_str(), binary );
 		oroFunction function;
 		oroModule module;
-		oroError ee = oroModuleLoadData( &module, binary.data() );
+		oroError ee = oroModuleLoadData( &module, binary.data(), binary.size() );
 		ee = oroModuleGetFunction( &function, module, "testKernel" );
 		int x = 123;
 		const void* args[] = { &x };

diff --git a/UnitTest/testKernel.cl b/UnitTest/testKernel.cl
@@ -0,0 +1,5 @@
+__kernel void testKernel( __global int* a )
+{
+	int tid = get_global_id( 0 );
+	atomic_add(a, tid);
+}