gaspersic/triton
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Driver: now loading the backend dynamically on Linux

Browse Source
This commit is contained in:
Philippe Tillet
2015-08-25 12:41:21 -04:00
parent 868df9dad2
commit 67a35a62bd
44 changed files with 11808 additions and 13703 deletions
Download Patch File Download Diff File Expand all files Collapse all files

34
lib/driver/program.cpp
View File

@@ -4,9 +4,7 @@
#include "isaac/driver/program.h"
#include "isaac/driver/context.h"
#ifdef ISAAC_WITH_CUDA
#include "helpers/cuda/vector.hpp"
#endif
#include "helpers/ocl/infos.hpp"
#include "sha1.hpp"
@@ -23,7 +21,6 @@ Program::Program(Context const & context, std::string const & source) : backend_
std::string cache_path = context.cache_path_;
switch(backend_)
{
#ifdef ISAAC_WITH_CUDA
case CUDA:
{
@@ -34,7 +31,7 @@ Program::Program(Context const & context, std::string const & source) : backend_
//Load cached program
if(cache_path.size() && std::ifstream(fname, std::ios::binary))
{
cuda::check(cuModuleLoad(&h_.cu(), fname.c_str()));
cuda::check(dispatch::cuModuleLoad(&h_.cu(), fname.c_str()));
break;
}
@@ -43,25 +40,25 @@ Program::Program(Context const & context, std::string const & source) : backend_
const char * includes[] = {"helper_math.h"};
const char * src[] = {helpers::cuda::vector};
nvrtc::check(nvrtcCreateProgram(&prog, source.c_str(), NULL, 1, src, includes));
nvrtc::check(dispatch::nvrtcCreateProgram(&prog, source.c_str(), NULL, 1, src, includes));
try{
const char * options[] = {"--gpu-architecture=compute_52", "--restrict"};
nvrtc::check(nvrtcCompileProgram(prog, 2, options));
nvrtc::check(dispatch::nvrtcCompileProgram(prog, 2, options));
}catch(nvrtc::exception::compilation const &)
{
size_t logsize;
nvrtc::check(nvrtcGetProgramLogSize(prog, &logsize));
nvrtc::check(dispatch::nvrtcGetProgramLogSize(prog, &logsize));
std::string log(logsize, 0);
nvrtc::check(nvrtcGetProgramLog(prog, (char*)log.data()));
nvrtc::check(dispatch::nvrtcGetProgramLog(prog, (char*)log.data()));
std::cout << "Compilation failed:" << std::endl;
std::cout << log << std::endl;
}
size_t ptx_size;
nvrtc::check(nvrtcGetPTXSize(prog, &ptx_size));
nvrtc::check(dispatch::nvrtcGetPTXSize(prog, &ptx_size));
std::vector<char> ptx(ptx_size);
nvrtc::check(nvrtcGetPTX(prog, ptx.data()));
cuda::check(cuModuleLoadDataEx(&h_.cu(), ptx.data(), 0, NULL, NULL));
nvrtc::check(dispatch::nvrtcGetPTX(prog, ptx.data()));
cuda::check(dispatch::cuModuleLoadDataEx(&h_.cu(), ptx.data(), 0, NULL, NULL));
//Save cached program
if (cache_path.size())
@@ -77,8 +74,8 @@ Program::Program(Context const & context, std::string const & source) : backend_
// oss.close();
// system(("/usr/local/cuda-7.0/bin/nvcc " + sha1 + ".cu -gencode arch=compute_50,code=sm_50 -cubin").c_str());
// system(("perl /home/philippe/Development/maxas/maxas.pl -e " + sha1 + ".cubin > " + sha1 + ".sass").c_str());
// system(("perl /home/philippe/Development/maxas/maxas.pl -i --noreuse" + sha1 + ".sass " + sha1 + ".cubin").c_str());
// system(("perl /maxas.pl -e " + sha1 + ".cubin > " + sha1 + ".sass").c_str());
// system(("perl /maxas.pl -i --noreuse" + sha1 + ".sass " + sha1 + ".cubin").c_str());
// std::ifstream ifs(sha1 + ".cubin");
// std::cout << sha1 << std::endl;
@@ -90,11 +87,10 @@ Program::Program(Context const & context, std::string const & source) : backend_
// str.assign((std::istreambuf_iterator<char>(ifs)),
// std::istreambuf_iterator<char>());
// cuda::check(cuModuleLoadDataEx(&h_.cu(), str.c_str(), 0, NULL, NULL));
// cuda::check(dispatch::cuModuleLoadDataEx(&h_.cu(), str.c_str(), 0, NULL, NULL));
break;
}
#endif
case OPENCL:
{
cl_int err;
@@ -118,18 +114,18 @@ Program::Program(Context const & context, std::string const & source) : backend_
buffer.resize(len);
cached.read((char*)buffer.data(), std::streamsize(len));
char* cbuffer = buffer.data();
h_.cl() = clCreateProgramWithBinary(context_.h_.cl(), static_cast<cl_uint>(devices.size()), devices.data(), &len, (const unsigned char **)&cbuffer, NULL, &err);
h_.cl() = dispatch::clCreateProgramWithBinary(context_.h_.cl(), static_cast<cl_uint>(devices.size()), devices.data(), &len, (const unsigned char **)&cbuffer, NULL, &err);
ocl::check(err);
ocl::check(clBuildProgram(h_.cl(), static_cast<cl_uint>(devices.size()), devices.data(), build_opt.c_str(), NULL, NULL));
ocl::check(dispatch::clBuildProgram(h_.cl(), static_cast<cl_uint>(devices.size()), devices.data(), build_opt.c_str(), NULL, NULL));
return;
}
}
std::size_t srclen = source.size();
const char * csrc = source.c_str();
h_.cl() = clCreateProgramWithSource(context_.h_.cl(), 1, &csrc, &srclen, &err);
h_.cl() = dispatch::clCreateProgramWithSource(context_.h_.cl(), 1, &csrc, &srclen, &err);
try{
ocl::check(clBuildProgram(h_.cl(), static_cast<cl_uint>(devices.size()), devices.data(), build_opt.c_str(), NULL, NULL));
ocl::check(dispatch::clBuildProgram(h_.cl(), static_cast<cl_uint>(devices.size()), devices.data(), build_opt.c_str(), NULL, NULL));
}catch(ocl::exception::build_program_failure const &){
for(std::vector<cl_device_id>::const_iterator it = devices.begin(); it != devices.end(); ++it)
{