From 43cb9020c48a681d0ebd5d48883b1d7ce8f25ec9 Mon Sep 17 00:00:00 2001
From: Andreas Suter <andreas.suter@psi.ch>
Date: Tue, 9 Jun 2020 12:55:55 +0200
Subject: [PATCH] adapted for CUDA 11

---
 CMakeLists.txt                   |  2 +-
 ReadMe.first                     |  2 +-
 src/CUDA/CudaChiSquareRuntime.cu | 42 +++++++++++++++++++++++++-------
 3 files changed, 35 insertions(+), 11 deletions(-)

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 0ea9bdb..502d883 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -2,7 +2,7 @@ CMAKE_MINIMUM_REQUIRED (VERSION 3.2)
 PROJECT (DKS)
 SET (DKS_VERSION_MAJOR 1)
 SET (DKS_VERSION_MINOR 1)
-SET (DKS_VERSION_PATCH 2)
+SET (DKS_VERSION_PATCH 3)
 set (DKS_VERSION ${DKS_VERSION_MAJOR}.${DKS_VERSION_MINOR}.${DKS_VERSION_PATCH})
 SET (PACKAGE \"dks\")
 SET (PACKAGE_BUGREPORT \"locans.uldis@psi.ch\")
diff --git a/ReadMe.first b/ReadMe.first
index 37686bd..871cd43 100644
--- a/ReadMe.first
+++ b/ReadMe.first
@@ -1,7 +1,7 @@
 ##################################################################
 #
 # Name:		Dynamic Kernel Scheduler
-# Version:	1.0
+# Version:	1.1
 # Author: 	Uldis Locans
 # Contacts:	locans.uldis@psi.ch
 #
diff --git a/src/CUDA/CudaChiSquareRuntime.cu b/src/CUDA/CudaChiSquareRuntime.cu
index ebbbcd0..6eba81d 100644
--- a/src/CUDA/CudaChiSquareRuntime.cu
+++ b/src/CUDA/CudaChiSquareRuntime.cu
@@ -86,15 +86,19 @@ int CudaChiSquareRuntime::compileProgram(std::string function, bool mlh) {
 
   //create program
   nvrtcProgram prog;
-  //std::cout << cudaProg.c_str() << std::endl;
-  nvrtcCreateProgram(&prog, cudaProg.c_str(), "chiSquareRuntime.cu", 0, NULL, NULL);
+//  std::cout << cudaProg.c_str() << std::endl;
+  nvrtcResult createResult = nvrtcCreateProgram(&prog, cudaProg.c_str(), "chiSquareRuntime.cu", 0, NULL, NULL);
+  if (createResult != NVRTC_SUCCESS) {
+    DEBUG_MSG("Program creation failed!");
+    return DKS_ERROR;
+  }
 
   //compile program
-  const char *opts[] = {"-fmad=false", ""};
-  int numopts = 1;
+  const char *opts[] = {"-arch=compute_35", "-fmad=false", ""};
+  int numopts = 2;
   if (mlh) {
-    opts[1] = "-DMLH";
-    numopts = 2;
+    opts[2] = "-DMLH";
+    numopts = 3;
   }
 
   nvrtcResult compileResults = nvrtcCompileProgram(prog, numopts, opts);
@@ -118,7 +122,11 @@ int CudaChiSquareRuntime::compileProgram(std::string function, bool mlh) {
   if (ptx_m != NULL)
     delete[] ptx_m;
   size_t ptxSize; 
-  nvrtcGetPTXSize(prog, &ptxSize); 
+  nvrtcResult ptxSizeResult = nvrtcGetPTXSize(prog, &ptxSize);
+  if (ptxSizeResult != NVRTC_SUCCESS) {
+    DEBUG_MSG("PTX get size error!");
+    return DKS_ERROR;
+  }
   ptx_m = new char[ptxSize]; 
   nvrtcResult nvrtcPTXResult = nvrtcGetPTX(prog, ptx_m);  
   
@@ -127,10 +135,26 @@ int CudaChiSquareRuntime::compileProgram(std::string function, bool mlh) {
     return DKS_ERROR;
   }
 
+  // add some additional diagnostics
+  const int buffer_size = 8192;
+  CUjit_option options[3];
+  void* values[3];
+  char error_log[buffer_size];
+  int err;
+  options[0] = CU_JIT_ERROR_LOG_BUFFER;
+  values[0]  = (void*)error_log;
+  options[1] = CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES;
+  values[1]  = (void*)buffer_size;
+  options[2] = CU_JIT_TARGET_FROM_CUCONTEXT;
+  values[2]  = 0;
   //load module from ptx
-  CUresult loadResult = cuModuleLoadDataEx(&module_m, ptx_m, 0, 0, 0); 
+  CUresult loadResult = cuModuleLoadDataEx(&module_m, ptx_m, 3, options, values); 
   if (loadResult != CUDA_SUCCESS) {
-    DEBUG_MSG("Load module from ptx failed!");
+    const char *err_msg;
+    cuGetErrorString(loadResult, &err_msg);
+    std::string msg = "Load module from ptx failed! (" + std::to_string(loadResult) + ") : " + err_msg;
+    DEBUG_MSG(msg);
+    DEBUG_MSG(error_log);
     return DKS_ERROR;
   }