From e941b2b761bfcddfcd27bb00271efb2baf2c9734 Mon Sep 17 00:00:00 2001
From: gregrodgers <grodgers@purdue.edu>
Date: Mon, 11 May 2015 13:45:57 -0500
Subject: [PATCH] Fix multiple CL files problem

---
 bin/snack.sh    |  4 ++--
 bin/snk_genw.sh | 25 ++++++++++++++++++++++---
 2 files changed, 24 insertions(+), 5 deletions(-)

diff --git a/bin/snack.sh b/bin/snack.sh
index 94e43d1..44427d5 100755
--- a/bin/snack.sh
+++ b/bin/snack.sh
@@ -496,7 +496,7 @@ else
 if [ $DRYRUN ] ; then
    echo "hexdump -v -e '""0x"" 1/1 ""%02X"" "",""' $BRIGDIR/$BRIGNAME "
 else
-   echo "char HSA_BrigMem[] = {" > $FULLBRIGHFILE
+   echo "char ${SYMBOLNAME}_HSA_BrigMem[] = {" > $FULLBRIGHFILE
    hexdump -v -e '"0x" 1/1 "%02X" ","' $BRIGDIR/$BRIGNAME >> $FULLBRIGHFILE
    rc=$?
    if [ $rc != 0 ] ; then 
@@ -504,7 +504,7 @@ else
       exit $rc
    fi
    echo "};" >> $FULLBRIGHFILE
-   echo "size_t HSA_BrigMemSz = sizeof(HSA_BrigMem);" >> $FULLBRIGHFILE
+   echo "size_t ${SYMBOLNAME}_HSA_BrigMemSz = sizeof(${SYMBOLNAME}_HSA_BrigMem);" >> $FULLBRIGHFILE
 fi
 
 
diff --git a/bin/snk_genw.sh b/bin/snk_genw.sh
index 0158799..874e163 100755
--- a/bin/snk_genw.sh
+++ b/bin/snk_genw.sh
@@ -330,7 +330,7 @@ static hsa_status_t get_kernarg_memory_region(hsa_region_t region, void* data) {
 /* Stream specific globals */
 hsa_queue_t* Stream_CommandQ[SNK_MAX_STREAMS];
 snk_task_t   SNK_Tasks[SNK_MAX_TASKS];
-int          SNK_NextTaskId = 0 ;
+static int          SNK_NextTaskId = 0 ;
 
 /* Context(cl file) specific globals */
 hsa_ext_module_t*                _CN__BrigModule;
@@ -370,7 +370,7 @@ status_t _CN__InitContext(){
     /* printf("The maximum queue size is %u.\n", (unsigned int) queue_size); */
 
     /* Load the BRIG binary.  */
-    _CN__BrigModule = (hsa_ext_module_t*) &HSA_BrigMem;
+    _CN__BrigModule = (hsa_ext_module_t*) &_CN__HSA_BrigMem;
 
     /* Create hsa program.  */
     memset(&_CN__HsaProgram,0,sizeof(hsa_ext_program_t));
@@ -452,6 +452,7 @@ function write_KernelStatics_template(){
 hsa_executable_symbol_t          _KN__Symbol;
 int                              _KN__FK = 0 ; 
 status_t                         _KN__init();
+status_t                         _KN__stop();
 uint64_t                         _KN__Kernel_Object;
 uint32_t                         _KN__Kernarg_Segment_Size; /* May not need to be global */
 uint32_t                         _KN__Group_Segment_Size;
@@ -474,7 +475,7 @@ extern status_t _KN__init(){
 
     /* Extract the symbol from the executable.  */
     /* printf("Kernel name _KN__: Looking for symbol %s\n","__OpenCL__KN__kernel"); */
-    err = hsa_executable_get_symbol(_CN__Executable, "", "&__OpenCL__KN__kernel", _CN__Agent , 0, &_KN__Symbol);
+    err = hsa_executable_get_symbol(_CN__Executable, NULL, "&__OpenCL__KN__kernel", _CN__Agent , 0, &_KN__Symbol);
     ErrorCheck(Extract the symbol from the executable, err);
 
     /* Extract dispatch information from the symbol */
@@ -492,6 +493,24 @@ extern status_t _KN__init(){
 
 } /* end of _KN__init */
 
+
+extern status_t _KN__stop(){
+    status_t err;
+    if (_CN__FC == 0 ) {
+       /* weird, but we cannot stop unless we initialized the context */
+       err = _CN__InitContext();
+       if ( err != STATUS_SUCCESS ) return err; 
+       _CN__FC = 1;
+    }
+    if ( _KN__FK == 1 ) {
+        /*  Currently nothing kernel specific must be recovered */
+       _KN__FK = 0;
+    }
+    return STATUS_SUCCESS;
+
+} /* end of _KN__stop */
+
+
 EOF
 }