fixed bug in CUDA parser with macros

parser/instrument.py

@@ -2,6 +2,7 @@
 import os
 import tempfile
 import sys
+import linecache
 from collections import defaultdict
 from tokenizer import Tokenizer
 from match import Match, FunctionType
@@ -42,7 +43,15 @@ def deprocess(self):
     dp = Deprocess(self.preFileName, tmpFname)
     if verbose(): print('Running de-processor...')
     dp.run()
-    if verbose(): print('... de-preprocessor done.')
+    if verbose(): 
+      print('... de-preprocessor done.')
+      print('Deprocessed file:', self.deprocessedFile)
+      with open(self.deprocessedFile, 'r') as fd:
+        i = 1
+        for l in fd:
+          print("{n:3d}: {line}".format(n=i, line=l[:-1]))
+          i += 1  
+
     #os.close(tmpFd)
     #if 'FPC_LEAVE_TEMP_FILES' not in os.environ:
     #  os.remove(tmpFname)
@@ -154,6 +163,34 @@ def produceInstrumentedLines(self):
   def is_omitted_line(self, file_name: str, line: int):
     return self.conf.is_line_omitted(file_name, line)
 
+  def isLineInDeviceCode(self, line: int):
+    for i in self.deviceDclLines:
+      line_begin  = i[0]
+      line_end    = i[1]
+      if line >= line_begin and line <= line_end:
+        return True
+    return False
+
+# Old version ----------------
+#  def instrument(self):
+#    fileName, ext = os.path.splitext(self.sourceFileName)
+#    self.instrumentedFileName = fileName+'_inst'+ext
+#    with open(self.sourceFileName, 'r') as fd:
+#      with open(self.instrumentedFileName, 'w') as outFile:
+#        l = 0
+#        for line in fd:
+#          l += 1
+#          if l in self.transformedLines.keys():
+#            if not self.is_omitted_line(self.sourceFileName, l):
+#              newLine = self.transformedLines[l]
+#              if verbose(): print(newLine[:-1])
+#              outFile.write(newLine[:-1]+'\n')
+#            else:
+#              outFile.write(line[:-1]+'\n')
+#          else:
+#            if verbose(): print(line[:-1])
+#            outFile.write(line[:-1]+'\n')
+
   def instrument(self):
     fileName, ext = os.path.splitext(self.sourceFileName)
     self.instrumentedFileName = fileName+'_inst'+ext
@@ -170,8 +207,13 @@ def instrument(self):
             else:
               outFile.write(line[:-1]+'\n')
           else:
-            if verbose(): print(line[:-1])
-            outFile.write(line[:-1]+'\n')
+            if self.isLineInDeviceCode(l):
+              newLine = linecache.getline(self.deprocessedFile, l)
+              outFile.write(newLine[:-1]+'\n')
+              if verbose(): print(newLine[:-1])
+            else:
+              if verbose(): print(line[:-1])
+              outFile.write(line[:-1]+'\n')
 
   def getInstrumentedFileName(self):
     return self.instrumentedFileName

parser/nvcc_fpchecker.py

@@ -25,7 +25,7 @@
 # --------------------------------------------------------------------------- #
 
 # File extensions that can have CUDA code
-CUDA_EXTENSION = ['.cu', '.cuda'] + ['.C', '.cc', '.cpp', '.CPP', '.c++', '.cp', '.cxx']
+CUDA_EXTENSION = ['.cu', '.cuda'] + ['.C', '.cc', '.cpp', '.CPP', '.c++', '.cp', '.cxx', '.c']
 
 # --------------------------------------------------------------------------- #
 # --- Classes --------------------------------------------------------------- #
@@ -191,6 +191,7 @@ def compileInstrumentedFile(self):
       cmd.instrumentSource()
       cmd.compileInstrumentedFile()
       logMessage('Instrumented: ' + cmd.instrumentedFile)
+      print("#FPCHECKER: Instrumented:", cmd.instrumentedFile)
     except Exception as e: # Fall back to original command
       if verbose():
         logMessage(str(e))

tests/parser/static/test_instrument_macro_unused_parms/test_instrument_macro_unused_parms.py

@@ -0,0 +1,105 @@
+import os
+import pathlib
+import sys
+import subprocess
+import pytest
+
+sys.path.insert(1, str(pathlib.Path(__file__).parent.absolute())+"/../../../../parser")
+#sys.path.insert(1, '/usr/workspace/wsa/laguna/fpchecker/FPChecker/parser')
+from tokenizer import Tokenizer
+from instrument import Instrument
+
+RUNTIME='../../../../src/Runtime_parser.h'
+
+prog_1 = """
+
+#define MULTI_LINE_MACRO(a, b, c, d, e) { a = b + c; }
+
+__device__ void comp(double *a, double b, double c) {
+  double tmp1 = 0.0;
+  double tmp2 = 0.1;
+
+  MULTI_LINE_MACRO( a[0],
+                    b, c,
+                    tmp1,
+                    tmp2 ); tmp1 = tmp2;
+}
+
+__host__ void comp2(int N) {
+
+#pragma omp parallel
+#pragma omp for
+for (int i=0; i<N; i++) {
+  // do something with i
+}
+
+}
+
+"""
+
+def setup_module(module):
+  THIS_DIR = os.path.dirname(os.path.abspath(__file__))
+  os.chdir(THIS_DIR)
+
+def teardown_module(module):
+  cmd = ["rm -f *.o *.ii *.cu"]
+  cmdOutput = subprocess.check_output(cmd, stderr=subprocess.STDOUT, shell=True)
+
+def preprocessFile(prog_name: str):
+  cmd = ['nvcc -E '+prog_name+'.cu -o '+prog_name+'.ii']
+  cmdOutput = subprocess.check_output(cmd, stderr=subprocess.STDOUT, shell=True)
+
+def createFile(prog: str, prog_name: str):
+  with open(prog_name+'.cu', 'w') as fd:
+    fd.write(prog)
+    fd.write('\n')
+  preprocessFile(prog_name)
+
+def instrument(prog_name: str):
+  pass
+  preFileName = prog_name+'.ii'
+  sourceFileName = prog_name+'.cu'
+  inst = Instrument(preFileName, sourceFileName)
+  inst.deprocess()
+  inst.findDeviceDeclarations()
+  inst.findAssigments()
+  inst.produceInstrumentedLines()
+  inst.instrument()
+
+def compileProggram(prog_name: str):
+  try:
+    cmd = ['nvcc -std=c++11 -c -include '+RUNTIME+' '+prog_name+'_inst.cu']
+    cmdOutput = subprocess.check_output(cmd, stderr=subprocess.STDOUT, shell=True)
+  except Exception as e:
+    print(e)
+    exit()
+
+def countInstrumentationCalls(prog_name: str):
+  ret = 0
+  with open(prog_name+'_inst.cu', 'r') as fd:
+    for l in fd.readlines():
+      for w in l.split():
+        if '_FPC_CHECK_' in w:
+          ret += 1
+  return ret
+
+def inst_program(prog: str, prog_name: str, num_inst: int):
+  try:
+    createFile(prog, prog_name)
+    instrument(prog_name)
+    compileProggram(prog_name)
+    n = countInstrumentationCalls(prog_name)
+    assert n == num_inst
+    return True
+  except Exception as e:
+    print(e)
+    return False
+
+#@pytest.mark.xfail(reason="Known parser issue with macros")
+def test_1():
+  os.environ['FPC_VERBOSE'] = '1'
+  inst_program(prog_1, 'prog_1', 1)
+
+if __name__ == '__main__':
+  test_1()
+

tests/parser/static/test_instrument_multiline_macro/test_instrument_multiline_macro.py

@@ -2,6 +2,7 @@
 import pathlib
 import sys
 import subprocess
+import pytest
 
 sys.path.insert(1, str(pathlib.Path(__file__).parent.absolute())+"/../../../../parser")
 #sys.path.insert(1, '/usr/workspace/wsa/laguna/fpchecker/FPChecker/parser')
@@ -75,8 +76,12 @@ def instrument(prog_name: str):
   inst.instrument()
 
 def compileProggram(prog_name: str):
-  cmd = ['nvcc -std=c++11 -c -include '+RUNTIME+' '+prog_name+'_inst.cu']
-  cmdOutput = subprocess.check_output(cmd, stderr=subprocess.STDOUT, shell=True)
+  try:
+    cmd = ['nvcc -std=c++11 -c -include '+RUNTIME+' '+prog_name+'_inst.cu']
+    cmdOutput = subprocess.check_output(cmd, stderr=subprocess.STDOUT, shell=True)
+  except Exception as e:
+    print(e)
+    exit()
 
 def countInstrumentationCalls(prog_name: str):
   ret = 0
@@ -93,12 +98,13 @@ def inst_program(prog: str, prog_name: str, num_inst: int):
     instrument(prog_name)
     compileProggram(prog_name)
     n = countInstrumentationCalls(prog_name)
-    #assert n == num_inst
+    assert n == num_inst
     return True
   except Exception as e:
     print(e)
     return False
 
+@pytest.mark.xfail(reason="Known parser issue with macros")
 def test_1():
   os.environ['FPC_VERBOSE'] = '1'
   inst_program(prog_1, 'prog_1', 1)

tests/parser/static/test_match_device_host_spaces/dot_product_inst.cu

@@ -1,10 +1,10 @@
 
 #include <stdio.h>
 
-__device__ void mul(double a, double b, double *res)
+__attribute__((device)) void mul(double a, double b, double *res)
 {
   *res = _FPC_CHECK_D_(a * b, 6, "/usr/WS1/laguna/fpchecker/FPChecker/tests/parser/static/test_match_device_host_spaces/dot_product.cu");
-  // NaN
+
   *res = _FPC_CHECK_D_((*res)-(*res) / (*res)-(*res), 8, "/usr/WS1/laguna/fpchecker/FPChecker/tests/parser/static/test_match_device_host_spaces/dot_product.cu");
 }
 
@@ -31,16 +31,16 @@ void calc(double *x, int s) {
   }
 }*/
 
-    __device__     __host__     void   comp(double *x, int s) {
+    __attribute__((device)) __attribute__((host)) void comp(double *x, int s) {
   for (int i=0; i < s; ++i) {
     x[i] = _FPC_CHECK_HD_(x[i] + 3.1, 36, "/usr/WS1/laguna/fpchecker/FPChecker/tests/parser/static/test_match_device_host_spaces/dot_product.cu");
   }
 }
 
-  __device__    
+  __attribute__((device))
 
 
- __host__     void   comp2(double *x, int s) {
+ __attribute__((host)) void comp2(double *x, int s) {
   for (int i=0; i < s; ++i) {
     x[i] = _FPC_CHECK_HD_(x[i] + 3.1, 45, "/usr/WS1/laguna/fpchecker/FPChecker/tests/parser/static/test_match_device_host_spaces/dot_product.cu");
   }