new nvcc wrapper

parser/colors.py

@@ -0,0 +1,10 @@
+def prGreen(skk):
+  print("\033[92m{}\033[00m" .format(skk))
+
+def prCyan(skk):
+  print("\033[96m{}\033[00m" .format(skk))
+
+def prRed(skk):
+  print("\033[91m{}\033[00m" .format(skk))
+
+

parser/instrument.py

@@ -22,6 +22,7 @@ def __init__(self, preprocessedFile, srcFile):
     self.PRE_DEVICE       = '_FPC_CHECK_D_'
     self.PRE_HOST_DEVICE  = '_FPC_CHECK_HD_'
     self.functionTypeMap = {} # key: token, value: function type
+    self.instrumentedFileName = None
 
   def __del__(self):
     if self.deprocessedFile:
@@ -118,8 +119,9 @@ def produceInstrumentedLines(self):
 
   def instrument(self):
     fileName, ext = os.path.splitext(self.sourceFileName)
+    self.instrumentedFileName = fileName+'_inst'+ext
     with open(self.sourceFileName, 'r') as fd:
-      with open(fileName+'_inst'+ext, 'w') as outFile:
+      with open(self.instrumentedFileName, 'w') as outFile:
         l = 0
         for line in fd:
           l += 1
@@ -131,6 +133,9 @@ def instrument(self):
             print(line[:-1])
             outFile.write(line[:-1]+'\n')
 
+  def getInstrumentedFileName(self):
+    return self.instrumentedFileName
+
 if __name__ == '__main__':
   preFileName = sys.argv[1]
   sourceFileName = sys.argv[2]

parser/nvcc_fpchecker.py

@@ -0,0 +1,183 @@
+#!/usr/bin/env python3
+
+import os
+import pathlib
+import subprocess
+import sys
+from colors import prGreen, prCyan, prRed
+from instrument import Instrument
+
+# --------------------------------------------------------------------------- #
+# --- Installation Paths ---------------------------------------------------- #
+# --------------------------------------------------------------------------- #
+
+# Main installation path
+FPCHECKER_PATH      = str(pathlib.Path(__file__).parent.absolute())
+FPCHECKER_LIB       = FPCHECKER_PATH+'/../lib/libfpchecker_plugin.so'
+FPCHECKER_RUNTIME   = FPCHECKER_PATH+'/../src/Runtime_parser.h'
+
+#
+# --------------------------------------------------------------------------- #
+# --- Global Variables ------------------------------------------------------ #
+# --------------------------------------------------------------------------- #
+
+# File extensions that can have CUDA code
+CUDA_EXTENSION = ['.cu', '.cuda'] + ['.C', '.cc', '.cpp', '.CPP', '.c++', '.cp', '.cxx']
+
+# --------------------------------------------------------------------------- #
+# --- Classes --------------------------------------------------------------- #
+# --------------------------------------------------------------------------- #
+
+class Command:
+  def __init__(self, cmd):
+    self.name = cmd[0]
+    self.parameters = cmd[1:]
+    self.preprocessedFile = None
+    self.instrumentedFile = None
+    self.outputFile = None
+
+  def executeOriginalCommand(self):
+    try:
+      cmd = ['nvcc'] + self.parameters
+      print('Executing original command:', cmd)
+      #cmdOutput = subprocess.call(cmd, stderr=subprocess.STDOUT, shell=True)
+      subprocess.run(' '.join(cmd), shell=True, check=True)
+    except subprocess.CalledProcessError as e:
+      prRed(e)
+      #sys.exit('FPCHECKER: nvcc error: ' + ' '.join(cmd))
+
+  def getOriginalCommand(self):
+    return ' '.join(['nvcc'] + self.parameters[1:])
+
+  # It it a compilation command?
+  def isCompileCommand(self) -> bool:
+    if ('-c' in self.parameters or
+        '--compile' in self.parameters or
+        '-dc' in self.parameters or
+        '--device-c' in self.parameters or
+        '-dw' in self.parameters or
+        '--device-w' in self.parameters):
+      return True
+    return False
+
+  # Is the command a link command?
+  def isLinkCommand(self) -> bool:
+    if ('-c' not in self.parameters and 
+        '--compile' not in self.parameters and
+        '-dw' not in self.parameters and
+        '--device-w' not in self.parameters and
+        '-cubin' not in self.parameters and
+        '-ptx' not in self.parameters and
+        '-fatbin' not in self.parameters and
+        ('-o' in self.parameters or '--output-file' in self.parameters)):
+      return True
+    return False
+
+  # Get the name of the cuda file to be compiled
+  # if the file exists.
+  def getCodeFileNameIfExists(self):
+    global CUDA_EXTENSION
+    fileName = None
+    for t in self.parameters:
+      for ext in CUDA_EXTENSION:
+        if t.endswith(ext):
+          fileName = t
+
+    if not fileName:
+      raise RuntimeError('FPCHECKER: Could not find source file') from None
+
+    return fileName
+
+  def getOutputFileIfExists(self):
+    for i in range(len(self.parameters)):
+      p = self.parameters[i]
+      if p == '-o' or p == '--output-file':
+        self.outputFile = self.parameters[i+1]
+        return self.parameters[i+1]
+    return None
+
+  # We transform the command and execute the pre-proecessor
+  def executePreprocessor(self):
+    source = self.getCodeFileNameIfExists()
+
+    # Copy the command parameters
+    newParams = self.parameters.copy()
+
+    outputFile = self.getOutputFileIfExists()
+    if outputFile:
+      self.preprocessedFile = outputFile + '.ii'
+      for i in range(len(newParams)):
+        p = self.parameters[i]
+        if p == '-o' or p == '--output-file':
+          newParams[i+1] = self.preprocessedFile
+          break
+    else:
+      self.preprocessedFile = source + '.ii'
+      newParams.append('-o')
+      newParams.append(self.preprocessedFile)
+
+    new_cmd = ['nvcc', '-E'] + newParams
+    try:
+      prGreen(new_cmd) 
+      cmdOutput = subprocess.run(' '.join(new_cmd), shell=True, check=True)
+    except subprocess.CalledProcessError as e:
+      prRed(e)
+      raise RuntimeError('FPCHECKER: Could not execute pre-processor') from e
+
+    return True
+
+  def instrumentSouce(self):
+    preFileName = self.preprocessedFile
+    sourceFileName = self.getCodeFileNameIfExists()
+    inst = Instrument(preFileName, sourceFileName)
+    inst.deprocess()
+    inst.findDeviceDeclarations()
+    print(inst.deviceDclLines)
+    inst.findAssigments()
+    inst.produceInstrumentedLines()
+    inst.instrument()
+    self.instrumentedFile = inst.getInstrumentedFileName()
+
+  def compileInstrumentedFile(self):
+    source = self.getCodeFileNameIfExists()
+    # Copy original command
+    new_cmd = ['nvcc', '-include', FPCHECKER_RUNTIME] + self.parameters
+    # Replace file by instrumented file
+    for i in range(len(new_cmd)):
+      p = new_cmd[i]
+      if p == source:
+        new_cmd[i] = self.instrumentedFile
+        break
+
+    # Change output file
+    if not self.outputFile:
+      fileName, ext = os.path.splitext(source)
+      newOutputFile = fileName + '.o'
+      new_cmd = new_cmd + ['-o', newOutputFile]
+
+    # Compile
+    try:
+      prGreen('Compiling: ' + ' '.join(new_cmd))
+      cmdOutput = subprocess.run(' '.join(new_cmd), shell=True, check=True)
+    except Exception as e:
+      prRed(e)
+      raise RuntimeError('FPCHECKER: Could not compile instrumented file') from e
+
+if __name__ == '__main__':
+  cmd = Command(sys.argv)
+
+  # Link command
+  if cmd.isLinkCommand():
+    #prGreen('Linking...')
+    cmd.executeOriginalCommand()
+  else:
+    # Compilation command
+    try:
+      cmd.executePreprocessor()
+      cmd.instrumentSouce()
+      cmd.compileInstrumentedFile()
+    except Exception as e:
+      # Fall back to original command
+      prRed(e)
+      cmd.executeOriginalCommand()
+

parser/test/Makefile

@@ -3,18 +3,23 @@
 #OP	= $(CLANG_PASS) -emit-llvm -O0 -x cuda --cuda-gpu-arch=sm_60 -g
 #LINK	= -L${CUDA_PATH}/lib64/ -lcuda -lcudart
 
+NVCC = /usr/workspace/wsa/laguna/fpchecker/FPChecker/parser/nvcc_fpchecker.py
+#NVCC = nvcc
+
 INCLUDE = -include /usr/workspace/wsa/laguna/fpchecker/FPChecker/src/Runtime_parser.h
 
 all:
 	cp -f dot_product.cu.orig dot_product.cu
 	cp -f main.cu.orig main.cu
-	nvcc  -arch=sm_70 -O0 -E dot_product.cu -o dot_product.cu.ii
-	nvcc  -arch=sm_70 -O0 -E main.cu -o main.cu.ii
-	python3 ../instrument.py dot_product.cu.ii dot_product.cu
-	python3 ../instrument.py main.cu.ii main.cu
-	nvcc $(INCLUDE) -arch=sm_70 -O0 -c dot_product_inst.cu 
-	nvcc $(INCLUDE) -arch=sm_70 -O0 -c main_inst.cu 
-	nvcc -o main main_inst.o dot_product_inst.o
+	#nvcc  -arch=sm_70 -O0 -E dot_product.cu -o dot_product.cu.ii
+	#nvcc  -arch=sm_70 -O0 -E main.cu -o main.cu.ii
+	#python3 ../instrument.py dot_product.cu.ii dot_product.cu
+	#python3 ../instrument.py main.cu.ii main.cu
+	#nvcc $(INCLUDE) -arch=sm_70 -O0 -c dot_product_inst.cu 
+	#nvcc $(INCLUDE) -arch=sm_70 -O0 -c main_inst.cu 
+	$(NVCC) --std=c++11 -arch=sm_70 -O0 -c dot_product.cu 
+	$(NVCC) --std=c++11 -arch=sm_70 -O0 -c main.cu 
+	$(NVCC) -o main main.o dot_product.o
 
 clean:
 	rm -rf *.cu *.o main __pycache__ *_inst.* *.inst *.ii

src/Runtime_parser.h

@@ -13,6 +13,7 @@
 #include <math.h>
 #include <stdint.h>
 #include <limits.h>
+#include <type_traits>
 
 /// *** Warning ***
 /// Changing this file: Runtime.h
@@ -86,7 +87,11 @@ __device__ static int _FPC_FP64_IS_FLUSH_TO_ZERO(double x, double y, double z, i
 template<typename T>
 __device__ 
 T _FPC_CHECK_D_(T t, int x, const char *str) {
-  return t;
+  if (std::is_floating_point<T>::value) {
+    return _FPC_CHECK_(t, x, str);
+  } else {
+    return t;
+  }
 }
 
 template<typename T>