| 314 | |
| 315 | template<typename T> |
| 316 | cl_kernel getKernel( |
| 317 | std::string funcName, cl_context ctx, cl_device_id dev, cl_command_queue q, |
| 318 | const nonstd::span<Node* const> full_nodes, |
| 319 | nonstd::span<Node_ids const> full_ids, nonstd::span<int const> output_ids, |
| 320 | nonstd::span<oneapi::AParam<T, sycl::access_mode::write> const> ap, |
| 321 | bool is_linear) { |
| 322 | std::string devName; |
| 323 | { |
| 324 | std::lock_guard<std::mutex> lock(device_name_map_mutex); |
| 325 | |
| 326 | auto devNameIt = device_name_map.find(dev); |
| 327 | if (devNameIt == device_name_map.end()) { |
| 328 | size_t devNameSz; |
| 329 | CL_CHECK( |
| 330 | clGetDeviceInfo(dev, CL_DEVICE_NAME, 0, nullptr, &devNameSz)); |
| 331 | string newDevName(devNameSz, '\0'); |
| 332 | CL_CHECK(clGetDeviceInfo(dev, CL_DEVICE_NAME, devNameSz, |
| 333 | newDevName.data(), nullptr)); |
| 334 | device_name_map[dev] = newDevName; |
| 335 | devName = newDevName; |
| 336 | } else { |
| 337 | devName = devNameIt->second; |
| 338 | } |
| 339 | } |
| 340 | |
| 341 | vector<cl_kernel> kernels(10); |
| 342 | bool kernel_found; |
| 343 | string kernelHash = funcName + devName; |
| 344 | { |
| 345 | std::lock_guard<std::mutex> lock(kernel_map_mutex); |
| 346 | kernel_found = !(kernel_map.find(kernelHash) == end(kernel_map)); |
| 347 | } |
| 348 | if (kernel_found) { |
| 349 | std::lock_guard<std::mutex> lock(kernel_map_mutex); |
| 350 | kernels[0] = kernel_map[kernelHash]; |
| 351 | } else { |
| 352 | string jitstr = arrayfire::opencl::getKernelString( |
| 353 | funcName, full_nodes, full_ids, output_ids, is_linear, false, false, |
| 354 | ap[0].dims[2] > 1); |
| 355 | |
| 356 | cl_int err; |
| 357 | vector<const char*> jitsources = { |
| 358 | {arrayfire::oneapi::opencl::KParam_hpp, |
| 359 | arrayfire::oneapi::opencl::jit_cl, jitstr.c_str()}}; |
| 360 | vector<size_t> jitsizes = {arrayfire::oneapi::opencl::KParam_hpp_len, |
| 361 | arrayfire::oneapi::opencl::jit_cl_len, |
| 362 | jitstr.size()}; |
| 363 | |
| 364 | cl_program prog = clCreateProgramWithSource( |
| 365 | ctx, jitsources.size(), jitsources.data(), jitsizes.data(), &err); |
| 366 | |
| 367 | std::string options = getActiveDeviceBaseBuildFlags(); |
| 368 | |
| 369 | CL_CHECK_BUILD( |
| 370 | clBuildProgram(prog, 1, &dev, options.c_str(), nullptr, nullptr)); |
| 371 | |
| 372 | cl_uint ret_kernels = 0; |
| 373 | CL_CHECK( |
nothing calls this directly
no test coverage detected