@@ -196,9 +196,13 @@ class OpenClContext : public ComputeContext {
196196
197197private:
198198 void updateVariatesPool ();
199- void releaseMem (cl_mem& m);
200- void releaseKernel (cl_kernel& k);
201- void releaseProgram (cl_program& p);
199+
200+ void runHealthChecks ();
201+ std::string runHealthCheckProgram (const std::string& source, const std::string& kernelName);
202+
203+ static void releaseMem (cl_mem& m);
204+ static void releaseKernel (cl_kernel& k);
205+ static void releaseProgram (cl_program& p);
202206
203207 enum class ComputeState { idle, createInput, createVariates, calc };
204208
@@ -367,6 +371,92 @@ void OpenClContext::releaseProgram(cl_program& p) {
367371 }
368372}
369373
374+ std::string OpenClContext::runHealthCheckProgram (const std::string& source, const std::string& kernelName) {
375+
376+ struct CleanUp {
377+ std::vector<cl_program> p;
378+ std::vector<cl_kernel> k;
379+ std::vector<cl_mem> m;
380+ ~CleanUp () {
381+ for (auto & pgm : p)
382+ OpenClContext::releaseProgram (pgm);
383+ for (auto & krn : k)
384+ OpenClContext::releaseKernel (krn);
385+ for (auto & mem : m)
386+ OpenClContext::releaseMem (mem);
387+ }
388+ } cleanup;
389+
390+ const char * programPtr = source.c_str ();
391+
392+ cl_int err;
393+
394+ cl_program program = clCreateProgramWithSource (context_, 1 , &programPtr, NULL , &err);
395+ if (err != CL_SUCCESS) {
396+ return errorText (err);
397+ }
398+ cleanup.p .push_back (program);
399+
400+ err = clBuildProgram (program, 1 , &device_, NULL , NULL , NULL );
401+ if (err != CL_SUCCESS) {
402+ return errorText (err);
403+ }
404+
405+ cl_kernel kernel = clCreateKernel (program, kernelName.c_str (), &err);
406+ if (err != CL_SUCCESS) {
407+ return errorText (err);
408+ }
409+ cleanup.k .push_back (kernel);
410+
411+ cl_mem resultBuffer = clCreateBuffer (context_, CL_MEM_READ_WRITE, sizeof (cl_ulong), NULL , &err);
412+ if (err != CL_SUCCESS) {
413+ return errorText (err);
414+ }
415+ cleanup.m .push_back (resultBuffer);
416+
417+ err = clSetKernelArg (kernel, 0 , sizeof (cl_mem), &resultBuffer);
418+
419+ cl_event runEvent;
420+ constexpr std::size_t sizeOne = 1 ;
421+ err = clEnqueueNDRangeKernel (queue_, kernel, 1 , NULL , &sizeOne, NULL , 0 , NULL , &runEvent);
422+ if (err != CL_SUCCESS) {
423+ return errorText (err);
424+ }
425+
426+ cl_ulong result;
427+ err = clEnqueueReadBuffer (queue_, resultBuffer, CL_TRUE, 0 , sizeof (cl_ulong), &result, 1 , &runEvent, NULL );
428+ if (err != CL_SUCCESS) {
429+ return errorText (err);
430+ }
431+
432+ return std::to_string (result);
433+ }
434+
435+ void OpenClContext::runHealthChecks () {
436+ deviceInfo_.push_back (std::make_pair (" host_sizeof(cl_uint)" , std::to_string (sizeof (cl_uint))));
437+ deviceInfo_.push_back (std::make_pair (" host_sizeof(cl_ulong)" , std::to_string (sizeof (cl_ulong))));
438+ deviceInfo_.push_back (std::make_pair (" host_sizeof(cl_float)" , std::to_string (sizeof (cl_float))));
439+ deviceInfo_.push_back (std::make_pair (" host_sizeof(cl_double)" , std::to_string (sizeof (cl_double))));
440+
441+ std::string kernelGetUintSize =
442+ " __kernel void ore_get_uint_size(__global ulong* result) { result[0] = sizeof(uint); }" ;
443+ std::string kernelGetUlongSize =
444+ " __kernel void ore_get_ulong_size(__global ulong* result) { result[0] = sizeof(ulong); }" ;
445+ std::string kernelGetFloatSize =
446+ " __kernel void ore_get_float_size(__global ulong* result) { result[0] = sizeof(float); }" ;
447+ std::string kernelGetDoubleSize =
448+ " __kernel void ore_get_double_size(__global ulong* result) { result[0] = sizeof(double); }" ;
449+
450+ deviceInfo_.push_back (
451+ std::make_pair (" device_sizeof(uint)" , runHealthCheckProgram (kernelGetUintSize, " ore_get_uint_size" )));
452+ deviceInfo_.push_back (
453+ std::make_pair (" device_sizeof(ulong)" , runHealthCheckProgram (kernelGetUlongSize, " ore_get_ulong_size" )));
454+ deviceInfo_.push_back (
455+ std::make_pair (" device_sizeof(float)" , runHealthCheckProgram (kernelGetFloatSize, " ore_get_float_size" )));
456+ deviceInfo_.push_back (
457+ std::make_pair (" device_sizeof(double)" , runHealthCheckProgram (kernelGetDoubleSize, " ore_get_double_size" )));
458+ }
459+
370460void OpenClContext::init () {
371461
372462 if (initialized_) {
@@ -395,6 +485,8 @@ void OpenClContext::init() {
395485 " OpenClContext::OpenClContext(): error during clCreateCommandQueue(): " << errorText (err));
396486
397487 initialized_ = true ;
488+
489+ runHealthChecks ();
398490}
399491
400492std::pair<std::size_t , bool > OpenClContext::initiateCalculation (const std::size_t n, const std::size_t id,
@@ -683,7 +775,14 @@ void OpenClContext::updateVariatesPool() {
683775 624 * (nVariates_ * size_[currentId_ - 1 ] / 624 + (nVariates_ * size_[currentId_ - 1 ] % 624 == 0 ? 0 : 1 ));
684776
685777 cl_int err;
778+
686779 cl_mem oldBuffer = variatesPool_;
780+ struct OldBufferReleaser {
781+ OldBufferReleaser (cl_mem b) : b(b) {}
782+ ~OldBufferReleaser () { OpenClContext::releaseMem (b); }
783+ cl_mem b;
784+ } oldBufferReleaser (oldBuffer);
785+
687786 variatesPool_ = clCreateBuffer (context_, CL_MEM_READ_WRITE, fpSize * alignedSize, NULL , &err);
688787 QL_REQUIRE (err == CL_SUCCESS, " OpenClContext::updateVariatesPool(): error creating variates buffer with size "
689788 << fpSize * alignedSize << " bytes: " << errorText (err));
@@ -735,11 +834,6 @@ void OpenClContext::updateVariatesPool() {
735834 if (!waitList.empty ())
736835 clWaitForEvents (waitList.size (), &waitList[0 ]);
737836
738- // release old buffer
739-
740- if (variatesPoolSize_ > 0 )
741- releaseMem (oldBuffer);
742-
743837 // update current variates pool size
744838
745839 QL_REQUIRE (currentPoolSize == alignedSize, " OpenClContext::updateVariatesPool(): internal error, currentPoolSize = "
0 commit comments