diff options
author | Cedric Nugteren <web@cedricnugteren.nl> | 2016-04-28 21:14:17 +0200 |
---|---|---|
committer | Cedric Nugteren <web@cedricnugteren.nl> | 2016-04-28 21:14:17 +0200 |
commit | d9b21d7f4920b115d3fe7f2e3cce1f89eb762c10 (patch) | |
tree | c9a44f189fce5f6fc2456604dfe4c9d3e951e4e0 /include/internal | |
parent | d7ddbdeb1f416f56bc469d16c051551207274703 (diff) |
Fixed the cache to store binaries instead of OpenCL programs
Diffstat (limited to 'include/internal')
-rw-r--r-- | include/internal/cache.h | 30 | ||||
-rw-r--r-- | include/internal/clpp11.h | 25 | ||||
-rw-r--r-- | include/internal/routine.h | 16 |
3 files changed, 46 insertions, 25 deletions
diff --git a/include/internal/cache.h b/include/internal/cache.h index 44fad68d..fa33b78f 100644 --- a/include/internal/cache.h +++ b/include/internal/cache.h @@ -24,9 +24,9 @@ namespace clblast { namespace cache { // ================================================================================================= -// The cache of compiled OpenCL programs, along with some meta-data -struct ProgramCache { - Program program; +// The cache of compiled OpenCL binaries, along with some meta-data +struct BinaryCache { + std::string binary; std::string device_name; Precision precision; std::string routine_name_; @@ -41,28 +41,28 @@ struct ProgramCache { }; // The actual cache, implemented as a vector of the above data-type, and its mutex -static std::vector<ProgramCache> program_cache_; -static std::mutex program_cache_mutex_; +static std::vector<BinaryCache> binary_cache_; +static std::mutex binary_cache_mutex_; // ================================================================================================= -// Stores the compiled program in the cache -void StoreProgramToCache(const Program& program, const std::string &device_name, - const Precision &precision, const std::string &routine_name); +// Stores the compiled binary in the cache +void StoreBinaryToCache(const std::string& binary, const std::string &device_name, + const Precision &precision, const std::string &routine_name); -// Queries the cache and retrieves a matching program. Assumes that the match is available, throws +// Queries the cache and retrieves a matching binary. Assumes that the match is available, throws // otherwise. -const Program& GetProgramFromCache(const std::string &device_name, const Precision &precision, - const std::string &routine_name); +const std::string& GetBinaryFromCache(const std::string &device_name, const Precision &precision, + const std::string &routine_name); // Queries the cache to see whether or not the compiled kernel is already there -bool ProgramIsInCache(const std::string &device_name, const Precision &precision, - const std::string &routine_name); +bool BinaryIsInCache(const std::string &device_name, const Precision &precision, + const std::string &routine_name); // ================================================================================================= -// Clears the cache of stored program binaries -StatusCode ClearCompiledProgramCache(); +// Clears the cache of stored binaries +StatusCode ClearCache(); // ================================================================================================= } // namespace cache diff --git a/include/internal/clpp11.h b/include/internal/clpp11.h index 543d423a..b865ab1e 100644 --- a/include/internal/clpp11.h +++ b/include/internal/clpp11.h @@ -283,7 +283,7 @@ class Program { public: // Note that there is no constructor based on the regular OpenCL data-type because of extra state - // Regular constructor with memory management + // Source-based constructor with memory management explicit Program(const Context &context, std::string source): program_(new cl_program, [](cl_program* p) { CheckError(clReleaseProgram(*p)); delete p; }), length_(source.length()), @@ -294,6 +294,22 @@ class Program { CheckError(status); } + // Binary-based constructor with memory management + explicit Program(const Device &device, const Context &context, const std::string& binary): + program_(new cl_program, [](cl_program* p) { CheckError(clReleaseProgram(*p)); delete p; }), + length_(binary.length()), + source_(binary), + source_ptr_(&source_[0]) { + auto status1 = CL_SUCCESS; + auto status2 = CL_SUCCESS; + const cl_device_id dev = device(); + *program_ = clCreateProgramWithBinary(context(), 1, &dev, &length_, + reinterpret_cast<const unsigned char**>(&source_ptr_), + &status1, &status2); + CheckError(status1); + CheckError(status2); + } + // Compiles the device program and returns whether or not there where any warnings/errors BuildStatus Build(const Device &device, std::vector<std::string> &options) { auto options_string = std::accumulate(options.begin(), options.end(), std::string{" "}); @@ -322,7 +338,7 @@ class Program { return result; } - // Retrieves an intermediate representation of the compiled program + // Retrieves a binary or an intermediate representation of the compiled program std::string GetIR() const { auto bytes = size_t{0}; CheckError(clGetProgramInfo(*program_, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &bytes, nullptr)); @@ -338,7 +354,7 @@ class Program { private: std::shared_ptr<cl_program> program_; size_t length_; - std::string source_; + std::string source_; // Note: the source can also be a binary or IR const char* source_ptr_; }; @@ -633,7 +649,8 @@ class Kernel { // Launches the kernel while waiting for other events CheckError(clEnqueueNDRangeKernel(queue(), *kernel_, static_cast<cl_uint>(global.size()), nullptr, global.data(), local.data(), - waitForEventsPlain.size(), waitForEventsPlain.data(), + static_cast<cl_uint>(waitForEventsPlain.size()), + waitForEventsPlain.data(), event)); } diff --git a/include/internal/routine.h b/include/internal/routine.h index 013769d8..32be6012 100644 --- a/include/internal/routine.h +++ b/include/internal/routine.h @@ -84,19 +84,23 @@ class Routine { const bool upper = false, const bool lower = false, const bool diagonal_imag_zero = false); - // Stores a newly compiled program into the cache - void StoreProgramToCache(const Program& program) const { - return cache::StoreProgramToCache(program, device_name_, precision_, routine_name_); + // Stores a newly compiled binary into the cache + void StoreBinaryToCache(const std::string& binary) const { + return cache::StoreBinaryToCache(binary, device_name_, precision_, routine_name_); } // Queries the cache and retrieve either a matching program or a boolean whether a match exists. // The first assumes that the program is available in the cache and will throw an exception // otherwise. - const Program& GetProgramFromCache() const { - return cache::GetProgramFromCache(device_name_, precision_, routine_name_); + Program GetProgramFromCache() const { + auto& binary = cache::GetBinaryFromCache(device_name_, precision_, routine_name_); + auto program = Program(device_, context_, binary); + auto options = std::vector<std::string>(); + program.Build(device_, options); + return program; } bool ProgramIsInCache() const { - return cache::ProgramIsInCache(device_name_, precision_, routine_name_); + return cache::BinaryIsInCache(device_name_, precision_, routine_name_); } // Non-static variable for the precision. Note that the same variable (but static) might exist in |