LMU/DKS
5
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

adapted for CUDA 11

Browse Source
This commit is contained in:
suter_a
2020-06-09 12:55:55 +02:00
parent 3d946f666b
commit 43cb9020c4
3 changed files with 35 additions and 11 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
CMakeLists.txt
View File

@ -2,7 +2,7 @@ CMAKE_MINIMUM_REQUIRED (VERSION 3.2)
PROJECT (DKS) PROJECT (DKS)
SET (DKS_VERSION_MAJOR 1) SET (DKS_VERSION_MAJOR 1)
SET (DKS_VERSION_MINOR 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 (DKS_VERSION ${DKS_VERSION_MAJOR}.${DKS_VERSION_MINOR}.${DKS_VERSION_PATCH})
SET (PACKAGE \"dks\") SET (PACKAGE \"dks\")
SET (PACKAGE_BUGREPORT \"locans.uldis@psi.ch\") SET (PACKAGE_BUGREPORT \"locans.uldis@psi.ch\")

2
ReadMe.first
View File

@ -1,7 +1,7 @@
################################################################## ##################################################################
# #
# Name: Dynamic Kernel Scheduler # Name: Dynamic Kernel Scheduler
# Version: 1.0 # Version: 1.1
# Author: Uldis Locans # Author: Uldis Locans
# Contacts: locans.uldis@psi.ch # Contacts: locans.uldis@psi.ch
# #

42
src/CUDA/CudaChiSquareRuntime.cu
View File

@ -86,15 +86,19 @@ int CudaChiSquareRuntime::compileProgram(std::string function, bool mlh) {
//create program //create program
nvrtcProgram prog; nvrtcProgram prog;
//std::cout << cudaProg.c_str() << std::endl; // std::cout << cudaProg.c_str() << std::endl;
nvrtcCreateProgram(&prog, cudaProg.c_str(), "chiSquareRuntime.cu", 0, NULL, NULL); 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 //compile program
const char *opts[] = {"-fmad=false", ""}; const char *opts[] = {"-arch=compute_35", "-fmad=false", ""};
int numopts = 1; int numopts = 2;
if (mlh) { if (mlh) {
opts[1] = "-DMLH"; opts[2] = "-DMLH";
numopts = 2; numopts = 3;
} }
nvrtcResult compileResults = nvrtcCompileProgram(prog, numopts, opts); nvrtcResult compileResults = nvrtcCompileProgram(prog, numopts, opts);
@ -118,7 +122,11 @@ int CudaChiSquareRuntime::compileProgram(std::string function, bool mlh) {
if (ptx_m != NULL) if (ptx_m != NULL)
delete[] ptx_m; delete[] ptx_m;
size_t ptxSize; 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]; ptx_m = new char[ptxSize];
nvrtcResult nvrtcPTXResult = nvrtcGetPTX(prog, ptx_m); nvrtcResult nvrtcPTXResult = nvrtcGetPTX(prog, ptx_m);
@ -127,10 +135,26 @@ int CudaChiSquareRuntime::compileProgram(std::string function, bool mlh) {
return DKS_ERROR; 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 //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) { 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; return DKS_ERROR;
} }