From cc35bc64d20d6216b691021ebcefdc04605e9d4c Mon Sep 17 00:00:00 2001
From: Scott Biersdorff <scottb@godzilla.nic.uoregon.edu>
Date: Wed, 29 Aug 2012 17:46:36 -0700
Subject: [PATCH] Support for MIC offloading. The process to get TAU working on
 both the host and device at the same time is a little messy--in part because
 of gaps in Intel's MIC software stack which should hopefully be resolved in
 future releases.

See README.mic_offload and examples/mic for the steps to instrument
offload applications.

This does not work with opari. Given the convoluted pomp instrumentation
method this would be a substatial amount of work to extend to MIC offloading.


Former-commit-id: 099cf96afddf8b94ae981c336c59ddc0756260f0
---
 README.mic_offload        | 18 ++++++++++++++++++
 examples/mic/Makefile     | 22 ++++++++++++++++++++++
 examples/mic/sample.c     | 38 ++++++++++++++++++++++++++++++++++++++
 include/Makefile          |  4 +---
 src/Profile/Profiler.cpp  |  1 -
 src/Profile/TauInit.cpp   |  5 +++--
 tools/src/tau_compiler.sh | 11 +++++------
 7 files changed, 87 insertions(+), 12 deletions(-)
 create mode 100644 README.mic_offload
 create mode 100644 examples/mic/Makefile
 create mode 100755 examples/mic/sample.c

diff --git a/README.mic_offload b/README.mic_offload
new file mode 100644
index 000000000..514c60428
--- /dev/null
+++ b/README.mic_offload
@@ -0,0 +1,18 @@
+Configuring TAU for MIC offload applications.
+
+TAU can track both the host and mic sides of a MIC offloaded application. 
+
+1. Configure TAU with both '-arch=mic_linux' and '-arch=x86_64'.
+
+2. Set TAU makefile to x86_64 Makefile. Set TAU option '-optMICOffload', compile
+and link application with the TAU compiler wrappers.
+
+4. Set TAU environment variable: TAU_MIC_OFFLOAD=1, and MIC_PROXY_FS_ROOT=`pwd`
+then run the application.
+
+5. Host profile will be written out as profile.0.0.* and MIC profiles will be
+written out as profile.0.<pid>.*.
+
+NOTE: MIC offloading does not work with opari/pomp. Because it obfuscates the
+instrumentation process the compiler can not tell which portions of the
+measurement library needs to be compiled for the MIC architecture.
diff --git a/examples/mic/Makefile b/examples/mic/Makefile
new file mode 100644
index 000000000..da170b8a3
--- /dev/null
+++ b/examples/mic/Makefile
@@ -0,0 +1,22 @@
+default: offload
+
+sample_native.o: sample.c
+	tau_cc.sh -openmp -mmic -c $< -o $@
+
+sample_host.o: sample.c
+	tau_cc.sh -openmp -no-offload -c $< -o $@
+
+sample.o: sample.c
+	tau_cc.sh -openmp -c $< -o $@
+
+native: sample_native.o
+	tau_cc.sh -openmp -mmic sample_native.o -o $@
+	
+host: sample_host.o
+	tau_cc.sh -openmp -no-offload sample_host.o -o $@
+
+offload: sample.o
+	tau_cc.sh -tau_options='-optVerbose -optMICOffload' -openmp sample.o -o $@
+
+clean:
+	rm -f *.o native host offload
diff --git a/examples/mic/sample.c b/examples/mic/sample.c
new file mode 100755
index 000000000..c36bbdd25
--- /dev/null
+++ b/examples/mic/sample.c
@@ -0,0 +1,38 @@
+#include <stdio.h>
+#include <math.h>
+
+#define ITERATIONS 1000000000
+
+__attribute__((target(mic))) float compute_pi(int start)
+{
+	 float pi = 0.0f;
+	 float p16 = 1;
+	 int i;
+#pragma omp parallel for private(i) reduction(+:pi)
+	 for (i=start; i<start+ITERATIONS/2; i++)
+	 {	
+	 		p16 = pow(16,i);
+			pi += 1.0/p16 * (4.0/(8*i + 1) - 2.0/(8*i + 4) -
+	          1.0/(8*i + 5) - 1.0/(8*i + 6));
+	 }
+	 return pi;
+}
+
+int main()
+{
+	float pi = 0.0f;
+
+//First offload half the computation to the mic card.
+#pragma offload target (mic)
+	{
+	 	pi = compute_pi(0);
+	}
+		 
+//Compute the second half on the host.
+	{
+		pi += compute_pi(ITERATIONS/2);
+	}
+	
+	printf("Approximation of PI: %.20f.\n", pi);
+		
+}
diff --git a/include/Makefile b/include/Makefile
index 048138295..923506a59 100644
--- a/include/Makefile
+++ b/include/Makefile
@@ -1174,10 +1174,8 @@ TAUCOMPILEROPTS=  -optPdtDir="$(PDTDIR)/${PDTARCHDIR}"\
 	-optTauCC="$(TAU_CCOMPILER)" \
 	-optTauIncludes="$(TAU_INCLUDE) $(TAU_MPI_INCLUDE) $(TAU_SHMEM_INCLUDE) $(OPARIINCDIR)" \
 	-optTauDefs="$(TAU_DEFS)" $(TAU_IBM_NOCOMPINST) \
-        -optTauCompile="$(TAU_INCLUDE) $(TAU_MPI_COMPILE_INCLUDE) $(TAU_DEFS) "\
         -optLinking="$(TAU_MPI_FLIBS) $(TAU_LIBS) $(TAU_LDFLAGS) $(TAU_CXXLIBS) $(TAU_SHMEM_LIBS) "\
-        -optSharedLinking="-L$(TAU_LIB_DIR)/shared$(TAU_CONFIG) $(TAU_MPI_FLIBS)  $(TAU_SHMEM_LIBS) $(TAU_EXLIBS) $(TAU_LDFLAGS) $(TAU_CXXLIBS)"\
-	$(TAU_COMPILER_EXTRA_OPTIONS) \
+        -optSharedLinking="-L$(TAU_LIB_DIR)/shared$(TAU_CONFIG) $(TAU_MPI_FLIBS)  $(TAU_SHMEM_LIBS) $(TAU_EXLIBS) $(TAU_LDFLAGS) $(TAU_CXXLIBS) $(TAU_COMPILER_EXTRA_OPTIONS)" \
 				-optMICOffloadLinking="$(MIC_OFFLOAD_LINKING)"\
 				-optMICOffloadSharedLinking="$(MIC_OFFLOAD_SHARED_LINKING)"\
         -optIncludeMemory="$(TAU_INCLUDE_MEMORY)"\
diff --git a/src/Profile/Profiler.cpp b/src/Profile/Profiler.cpp
index 041249d1b..c29875622 100644
--- a/src/Profile/Profiler.cpp
+++ b/src/Profile/Profiler.cpp
@@ -257,7 +257,6 @@ void Profiler::Start(int tid) {
   RtsLayer::getUSecD(tid, StartTime);	  
   TimeStamp = (x_uint64) StartTime[0]; // USE COUNTER1 for tracing
 
-  
   /********************************************************************************/
   /*** Extras ***/
   /********************************************************************************/
diff --git a/src/Profile/TauInit.cpp b/src/Profile/TauInit.cpp
index f366f4802..9cd3fbe1e 100644
--- a/src/Profile/TauInit.cpp
+++ b/src/Profile/TauInit.cpp
@@ -481,6 +481,9 @@ extern "C" int Tau_init_initializeTAU() {
 	Tau_metadata_fillMetaData();
 #endif
 
+  tau_initialized = 1;
+  Tau_global_decr_insideTAU();
+
 #ifdef __MIC__
 	if (TauEnv_get_mic_offload())
 	{
@@ -489,8 +492,6 @@ extern "C" int Tau_init_initializeTAU() {
 	}
 #endif
 
-  tau_initialized = 1;
-  Tau_global_decr_insideTAU();
   return 0;
 }
 
diff --git a/tools/src/tau_compiler.sh b/tools/src/tau_compiler.sh
index 36c446f3b..1f10abf9c 100755
--- a/tools/src/tau_compiler.sh
+++ b/tools/src/tau_compiler.sh
@@ -773,11 +773,6 @@ for arg in "$@" ; do
 			optFixHashIf=$TRUE
 			echoIfDebug "\tFixing Hash-Ifs"
 			;;
-		    -opt*)
-			#Assume any other options should be passed on to the compiler.
-			argsRemaining="$argsRemaining ${arg%% *}"
-			;;
-
 		    -optMICOffloadLinking*)
 			optMICOffloadLinking="${arg#"-optMICOffloadLinking="} $optMICOffloadLinking"
 			echoIfDebug "\tLinking Options are: $optMICOffloadLinking"
@@ -790,6 +785,10 @@ for arg in "$@" ; do
 			optMICOffload=$TRUE
 			echoIfDebug "\tLinking for MIC Offloading"
 			;;
+		    -opt*)
+			#Assume any other options should be passed on to the compiler.
+			argsRemaining="$argsRemaining ${arg%% *}"
+			;;
 
 		esac #end case for parsing script Options
 		;;
@@ -1237,7 +1236,7 @@ if [ $optMICOffload == $TRUE ]; then
 	#	exit 1
 	#fi
 	#hybridLinking="$optLinking -offload-build -offload-ldopts='$optMICLinking'"
-	#echoIfDebug "Hybrid linking options: $hybridLinking"
+	echoIfDebug "MIC offload linking enabled."
 	if [ $optShared == $TRUE ]; then
 		optLinking=$optMICOffloadSharedLinking
 	else