Merge pull request #50 from amdadvtech/feature/ORO-0-only-bundle

Feature/oro 0 only bundle
GPUOpen-LibrariesAndSDKs · Mar 7, 2023 · c17bf53 · c17bf53
2 parents 7c9ad91 + 94ff379
commit c17bf53
Show file tree

Hide file tree

Showing 48 changed files with 36 additions and 80 deletions.
diff --git a/Orochi/Orochi.cpp b/Orochi/Orochi.cpp
@@ -338,6 +338,9 @@ oroError OROAPI oroGetDeviceProperties(oroDeviceProp* props, oroDevice dev)
 		e = cuDeviceGetAttribute( &props->computeMode, CU_DEVICE_ATTRIBUTE_COMPUTE_MODE, deviceId );
 		e = cuDeviceGetAttribute( &props->concurrentKernels, CU_DEVICE_ATTRIBUTE_CONCURRENT_KERNELS, deviceId );
 		e = cuDeviceGetAttribute( &props->ECCEnabled, CU_DEVICE_ATTRIBUTE_ECC_ENABLED, deviceId );
+		e = cuDeviceGetAttribute( &props->major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, deviceId );
+		e = cuDeviceGetAttribute( &props->minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, deviceId );
+
 		return oroSuccess;
 	}
 	return oroErrorUnknown;

diff --git a/UnitTest/bitcodes/generate_bitcodes.bat b/UnitTest/bitcodes/generate_bitcodes.bat
@@ -0,0 +1,2 @@
+call hipcc --cuda-device-only --offload-arch=gfx1030 --offload-arch=gfx1031 --offload-arch=gfx1032 --offload-arch=gfx1033 --offload-arch=gfx1034 --offload-arch=gfx1035 --offload-arch=gfx1036 --offload-arch=gfx1010 --offload-arch=gfx1011 --offload-arch=gfx1012 --offload-arch=gfx1013 --offload-arch=gfx900 --offload-arch=gfx906 -fgpu-rdc -c --gpu-bundle-output -c -emit-llvm ../moduleTestKernel.cpp
+call hipcc --cuda-device-only --offload-arch=gfx1030 --offload-arch=gfx1031 --offload-arch=gfx1032 --offload-arch=gfx1033 --offload-arch=gfx1034 --offload-arch=gfx1035 --offload-arch=gfx1036 --offload-arch=gfx1010 --offload-arch=gfx1011 --offload-arch=gfx1012 --offload-arch=gfx1013 --offload-arch=gfx900 --offload-arch=gfx906 -fgpu-rdc -c --gpu-bundle-output -c -emit-llvm ../moduleTestFunc.cpp
diff --git a/UnitTest/bitcodes/generate_bitcodes.sh b/UnitTest/bitcodes/generate_bitcodes.sh
@@ -1,30 +1,2 @@
 hipcc --cuda-device-only --offload-arch=gfx1030 --offload-arch=gfx1031 --offload-arch=gfx1032 --offload-arch=gfx1033 --offload-arch=gfx1034 --offload-arch=gfx1035 --offload-arch=gfx1036 --offload-arch=gfx1010 --offload-arch=gfx1011 --offload-arch=gfx1012 --offload-arch=gfx1013 --offload-arch=gfx900 --offload-arch=gfx906 -fgpu-rdc -c --gpu-bundle-output -c -emit-llvm ../moduleTestKernel.cpp
-
 hipcc --cuda-device-only --offload-arch=gfx1030 --offload-arch=gfx1031 --offload-arch=gfx1032 --offload-arch=gfx1033 --offload-arch=gfx1034 --offload-arch=gfx1035 --offload-arch=gfx1036 --offload-arch=gfx1010 --offload-arch=gfx1011 --offload-arch=gfx1012 --offload-arch=gfx1013 --offload-arch=gfx900 --offload-arch=gfx906 -fgpu-rdc -c --gpu-bundle-output -c -emit-llvm ../moduleTestFunc.cpp
-hipcc --offload-arch=gfx1030 ../moduleTestFunc.cpp -c -fgpu-rdc --cuda-device-only
-hipcc --offload-arch=gfx1031 ../moduleTestFunc.cpp -c -fgpu-rdc --cuda-device-only
-hipcc --offload-arch=gfx1032 ../moduleTestFunc.cpp -c -fgpu-rdc --cuda-device-only
-hipcc --offload-arch=gfx1033 ../moduleTestFunc.cpp -c -fgpu-rdc --cuda-device-only
-hipcc --offload-arch=gfx1034 ../moduleTestFunc.cpp -c -fgpu-rdc --cuda-device-only
-hipcc --offload-arch=gfx1035 ../moduleTestFunc.cpp -c -fgpu-rdc --cuda-device-only
-hipcc --offload-arch=gfx1036 ../moduleTestFunc.cpp -c -fgpu-rdc --cuda-device-only
-hipcc --offload-arch=gfx1010 ../moduleTestFunc.cpp -c -fgpu-rdc --cuda-device-only
-hipcc --offload-arch=gfx1011 ../moduleTestFunc.cpp -c -fgpu-rdc --cuda-device-only
-hipcc --offload-arch=gfx1012 ../moduleTestFunc.cpp -c -fgpu-rdc --cuda-device-only
-hipcc --offload-arch=gfx1013 ../moduleTestFunc.cpp -c -fgpu-rdc --cuda-device-only
-hipcc --offload-arch=gfx900 ../moduleTestFunc.cpp -c -fgpu-rdc --cuda-device-only
-hipcc --offload-arch=gfx906 ../moduleTestFunc.cpp -c -fgpu-rdc --cuda-device-only
-
-hipcc --offload-arch=gfx1030 ../moduleTestKernel.cpp -c -fgpu-rdc --cuda-device-only
-hipcc --offload-arch=gfx1031 ../moduleTestKernel.cpp -c -fgpu-rdc --cuda-device-only
-hipcc --offload-arch=gfx1032 ../moduleTestKernel.cpp -c -fgpu-rdc --cuda-device-only
-hipcc --offload-arch=gfx1033 ../moduleTestKernel.cpp -c -fgpu-rdc --cuda-device-only
-hipcc --offload-arch=gfx1034 ../moduleTestKernel.cpp -c -fgpu-rdc --cuda-device-only
-hipcc --offload-arch=gfx1035 ../moduleTestKernel.cpp -c -fgpu-rdc --cuda-device-only
-hipcc --offload-arch=gfx1036 ../moduleTestKernel.cpp -c -fgpu-rdc --cuda-device-only
-hipcc --offload-arch=gfx1010 ../moduleTestKernel.cpp -c -fgpu-rdc --cuda-device-only
-hipcc --offload-arch=gfx1011 ../moduleTestKernel.cpp -c -fgpu-rdc --cuda-device-only
-hipcc --offload-arch=gfx1012 ../moduleTestKernel.cpp -c -fgpu-rdc --cuda-device-only
-hipcc --offload-arch=gfx1013 ../moduleTestKernel.cpp -c -fgpu-rdc --cuda-device-only
-hipcc --offload-arch=gfx900 ../moduleTestKernel.cpp -c -fgpu-rdc --cuda-device-only
-hipcc --offload-arch=gfx906 ../moduleTestKernel.cpp -c -fgpu-rdc --cuda-device-only
diff --git a/UnitTest/bitcodes/generate_bitcodes_gfx1100.bat b/UnitTest/bitcodes/generate_bitcodes_gfx1100.bat
diff --git a/UnitTest/bitcodes/generate_bitcodes_gfx1100.sh b/UnitTest/bitcodes/generate_bitcodes_gfx1100.sh
diff --git a/UnitTest/bitcodes/generate_bitcodes_gfx1102.bat b/UnitTest/bitcodes/generate_bitcodes_gfx1102.bat
diff --git a/UnitTest/bitcodes/generate_bitcodes_gfx1102.sh b/UnitTest/bitcodes/generate_bitcodes_gfx1102.sh
diff --git a/UnitTest/bitcodes/generate_bitcodes_nvidia.bat b/UnitTest/bitcodes/generate_bitcodes_nvidia.bat
@@ -0,0 +1,2 @@
+nvcc -x cu -fatbin --device-c -arch=all ../moduleTestFunc.cpp
+nvcc -x cu -fatbin --device-c -arch=all ../moduleTestKernel.cpp
diff --git a/UnitTest/bitcodes/generate_bitcodes_nvidia.sh b/UnitTest/bitcodes/generate_bitcodes_nvidia.sh
@@ -1,4 +1,2 @@
-nvcc -arch=compute_80 -code="sm_80,sm_86,sm_87" -fatbin --device-c ../moduleTestFunc.cu
-nvcc -arch=compute_80 -code="sm_80,sm_86,sm_87" -fatbin --device-c ../moduleTestKernel.cu
-nvcc -cubin --device-c -arch=sm_80 ../moduleTestFunc.cu
-nvcc -cubin --device-c -arch=sm_80 ../moduleTestKernel.cu
+nvcc -fatbin --device-c -arch=all ../moduleTestFunc.cu
+nvcc -fatbin --device-c -arch=all ../moduleTestKernel.cu
diff --git a/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx1010.bc b/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx1010.bc
diff --git a/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx1011.bc b/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx1011.bc
diff --git a/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx1012.bc b/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx1012.bc
diff --git a/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx1013.bc b/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx1013.bc
diff --git a/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx1030.bc b/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx1030.bc
diff --git a/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx1031.bc b/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx1031.bc
diff --git a/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx1032.bc b/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx1032.bc
diff --git a/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx1033.bc b/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx1033.bc
diff --git a/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx1034.bc b/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx1034.bc
diff --git a/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx1035.bc b/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx1035.bc
diff --git a/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx1036.bc b/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx1036.bc
diff --git a/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx900.bc b/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx900.bc
diff --git a/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx906.bc b/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa-gfx906.bc
diff --git a/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa.bc b/UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa.bc
diff --git a/UnitTest/bitcodes/moduleTestFunc.cubin b/UnitTest/bitcodes/moduleTestFunc.cubin
diff --git a/UnitTest/bitcodes/moduleTestFunc.fatbin b/UnitTest/bitcodes/moduleTestFunc.fatbin
diff --git a/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx1010.bc b/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx1010.bc
diff --git a/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx1011.bc b/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx1011.bc
diff --git a/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx1012.bc b/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx1012.bc
diff --git a/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx1013.bc b/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx1013.bc
diff --git a/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx1030.bc b/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx1030.bc
diff --git a/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx1031.bc b/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx1031.bc
diff --git a/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx1032.bc b/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx1032.bc
diff --git a/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx1033.bc b/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx1033.bc
diff --git a/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx1034.bc b/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx1034.bc
diff --git a/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx1035.bc b/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx1035.bc
diff --git a/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx1036.bc b/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx1036.bc
diff --git a/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx900.bc b/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx900.bc
diff --git a/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx906.bc b/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa-gfx906.bc
diff --git a/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa.bc b/UnitTest/bitcodes/moduleTestKernel-hip-amdgcn-amd-amdhsa.bc
diff --git a/UnitTest/bitcodes/moduleTestKernel.cubin b/UnitTest/bitcodes/moduleTestKernel.cubin
diff --git a/UnitTest/bitcodes/moduleTestKernel.fatbin b/UnitTest/bitcodes/moduleTestKernel.fatbin
diff --git a/UnitTest/main.cpp b/UnitTest/main.cpp
@@ -149,7 +149,7 @@ void loadFile( const char* path, std::vector<char>& dst )
 		f.close();
 	}
 }
-
+#if 0
 TEST_F( OroTestBase, linkBc )
 {
 	oroDeviceProp props;
@@ -228,7 +228,7 @@ TEST_F( OroTestBase, linkBc )
 		ORORTCCHECK( oroModuleUnload( module ) );
 	}
 }
-
+#endif
 TEST_F( OroTestBase, link ) 
 {
 	oroDeviceProp props;
@@ -237,8 +237,10 @@ TEST_F( OroTestBase, link )
 	std::vector<char> data1;
 	const bool isAmd = oroGetCurAPI( 0 ) == ORO_API_HIP;
 
+	std::string arch = "-arch=sm_" + std::to_string( props.major ) + std::string( "0" );
+
 	std::vector<const char*> opts = isAmd ? std::vector<const char *>({ "-fgpu-rdc", "-c", "--cuda-device-only" })
-											:  std::vector<const char *>({ "--device-c", "-arch=sm_80" });
+											:  std::vector<const char *>({ "--device-c", arch.c_str() });
 	{
 		std::string code;
 		OrochiUtils::readSourceCode( "../UnitTest/moduleTestKernel.h", code );
@@ -310,7 +312,7 @@ TEST_F( OroTestBase, link )
 		ORORTCCHECK( oroModuleUnload( module ) );
 	}
 }
-
+#if 0
 TEST_F( OroTestBase, link_addFile )
 {
 	oroDeviceProp props;
@@ -385,6 +387,7 @@ TEST_F( OroTestBase, link_addFile )
 		ORORTCCHECK( oroModuleUnload( module ) );
 	}
 }
+#endif
 
 TEST_F( OroTestBase, link_null_name ) 
 {
@@ -393,9 +396,9 @@ TEST_F( OroTestBase, link_null_name )
 	std::vector<char> data0;
 	std::vector<char> data1;
 	const bool isAmd = oroGetCurAPI( 0 ) == ORO_API_HIP;
-
+	std::string arch = "-arch=sm_" + std::to_string( props.major ) + std::string( "0" );
 	std::vector<const char*> opts = isAmd ? std::vector<const char *>({ "-fgpu-rdc", "-c", "--cuda-device-only" })
-											:  std::vector<const char *>({ "--device-c", "-arch=sm_80" });
+											:  std::vector<const char *>({ "--device-c", arch.c_str() });
 	{
 		std::string code;
 		OrochiUtils::readSourceCode( "../UnitTest/moduleTestKernel.h", code );
@@ -526,14 +529,14 @@ TEST_F( OroTestBase, link_bundledBc_with_bc )
 	std::vector<char> data0;
 	std::vector<char> data1;
 	const bool isAmd = oroGetCurAPI( 0 ) == ORO_API_HIP;
-
+	std::string arch = "-arch=sm_" + std::to_string( props.major ) + std::string( "0" );
 	{
 		std::string bcFile = isAmd ? "../UnitTest/bitcodes/moduleTestFunc-hip-amdgcn-amd-amdhsa.bc" : "../UnitTest/bitcodes/moduleTestFunc.fatbin";
 		loadFile( bcFile.c_str(), data1 );
 	}
 	{
 		std::vector<const char*> opts = isAmd ? std::vector<const char *>({ "-fgpu-rdc", "-c", "--cuda-device-only" })
-											:  std::vector<const char *>({ "--device-c", "-arch=sm_80" });
+											:  std::vector<const char *>({ "--device-c", arch.c_str() });
 		std::string code;
 		OrochiUtils::readSourceCode( "../UnitTest/moduleTestKernel.h", code );
 		OrochiUtils::getData( m_device, code.c_str(), "../UnitTest/moduleTestKernel.h", &opts, data0 );
@@ -607,6 +610,7 @@ TEST_F( OroTestBase, link_bundledBc_with_bc_loweredName )
 	std::vector<char> data0;
 	std::vector<char> data1;
 	const bool isAmd = oroGetCurAPI( 0 ) == ORO_API_HIP;
+	std::string arch = "-arch=sm_" + std::to_string( props.major ) + std::string( "0" );
 	const char* funcName = "testKernel<0>";
 	std::string loweredNameStr;
 	orortcProgram prog;
@@ -616,7 +620,8 @@ TEST_F( OroTestBase, link_bundledBc_with_bc_loweredName )
 		loadFile( bcFile.c_str(), data1 );
 	}
 	{
-		std::vector<const char*> opts = isAmd ? std::vector<const char*>( { "-fgpu-rdc", "-c", "--cuda-device-only" } ) : std::vector<const char*>( { "--device-c", "-arch=sm_80" } );
+		std::vector<const char*> opts = isAmd ? std::vector<const char*>( { "-fgpu-rdc", "-c", "--cuda-device-only" } ) 
+			: std::vector<const char*>( { "--device-c", arch.c_str() } );
 		std::string code;
 
 		OrochiUtils::readSourceCode( "../UnitTest/moduleTestKernel_loweredName.h", code );

diff --git a/UnitTest/moduleTestFunc.cpp b/UnitTest/moduleTestFunc.cpp
@@ -1,4 +1,6 @@
+#if !defined( __CUDACC__ )
 #include <hip/hip_runtime.h>
+#endif
 
  __device__ void setInfo( int *x )
 {

diff --git a/UnitTest/moduleTestFunc.cu b/UnitTest/moduleTestFunc.cu
diff --git a/UnitTest/moduleTestKernel.cu b/UnitTest/moduleTestKernel.cu
diff --git a/UnitTest/premake5.lua b/UnitTest/premake5.lua
@@ -19,3 +19,7 @@ project "Unittest"
       files { "../contrib/gtest-1.6.0/gtest-all.cc" }
       sysincludedirs{ "../contrib/gtest-1.6.0/" }
       defines { "GTEST_HAS_TR1_TUPLE=0" }
+      if _OPTIONS["kernelcompile"] then
+        os.execute( "cd ./bitcodes/ && generate_bitcodes.bat" )
+        os.execute( "cd ./bitcodes/ && generate_bitcodes_nvidia.bat" )
+      end
diff --git a/contrib/hipew/src/hipew.cpp b/contrib/hipew/src/hipew.cpp
@@ -480,8 +480,8 @@ void hipewInit( int* resultDriver, int* resultRtc, hipuint32_t flags )
     _LIBRARY_FIND_CHECKED( rtcLib, hiprtcGetProgramLog );
     _LIBRARY_FIND_CHECKED( rtcLib, hiprtcGetProgramLogSize );
     _LIBRARY_FIND_CHECKED( rtcLib, hiprtcGetCode );
-    // _LIBRARY_FIND_CHECKED( rtcLib, hiprtcGetBitcodeSize );
-    // _LIBRARY_FIND_CHECKED( rtcLib, hiprtcGetBitcode );
+    _LIBRARY_FIND_CHECKED( rtcLib, hiprtcGetBitcodeSize );
+    _LIBRARY_FIND_CHECKED( rtcLib, hiprtcGetBitcode );
     _LIBRARY_FIND_CHECKED( rtcLib, hiprtcGetCodeSize );
     _LIBRARY_FIND_CHECKED( rtcLib, hiprtcLinkCreate );
     _LIBRARY_FIND_CHECKED( rtcLib, hiprtcLinkAddFile );

diff --git a/premake5.lua b/premake5.lua
@@ -8,6 +8,11 @@ newoption {
    description = "Use precompiled kernels"
 }
 
+newoption {
+   trigger = "kernelcompile",
+   description = "Compile kernels used for unit test"
+}
+
 function copydir(src_dir, dst_dir, filter, single_dst_dir)
 	if not os.isdir(src_dir) then
 		printError("'%s' is not an existing directory!", src_dir)
Original file line number	Diff line number	Diff line change
		@@ -0,0 +1,2 @@
		call hipcc --cuda-device-only --offload-arch=gfx1030 --offload-arch=gfx1031 --offload-arch=gfx1032 --offload-arch=gfx1033 --offload-arch=gfx1034 --offload-arch=gfx1035 --offload-arch=gfx1036 --offload-arch=gfx1010 --offload-arch=gfx1011 --offload-arch=gfx1012 --offload-arch=gfx1013 --offload-arch=gfx900 --offload-arch=gfx906 -fgpu-rdc -c --gpu-bundle-output -c -emit-llvm ../moduleTestKernel.cpp
		call hipcc --cuda-device-only --offload-arch=gfx1030 --offload-arch=gfx1031 --offload-arch=gfx1032 --offload-arch=gfx1033 --offload-arch=gfx1034 --offload-arch=gfx1035 --offload-arch=gfx1036 --offload-arch=gfx1010 --offload-arch=gfx1011 --offload-arch=gfx1012 --offload-arch=gfx1013 --offload-arch=gfx900 --offload-arch=gfx906 -fgpu-rdc -c --gpu-bundle-output -c -emit-llvm ../moduleTestFunc.cpp
Original file line number	Diff line number	Diff line change
		@@ -0,0 +1,2 @@
		nvcc -x cu -fatbin --device-c -arch=all ../moduleTestFunc.cpp
		nvcc -x cu -fatbin --device-c -arch=all ../moduleTestKernel.cpp