From 22fcaa07d6fd5dc61fade932dcf754180dd2dca7 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Mon, 24 Feb 2025 11:23:54 +0100 Subject: [PATCH 1/3] GPU: Do not try use OpenCL platforms where device query fails or which have 0 devices --- .../Base/opencl/GPUReconstructionOCL.cxx | 101 +++++++++--------- .../Base/opencl/GPUReconstructionOCL.h | 1 - 2 files changed, 51 insertions(+), 51 deletions(-) diff --git a/GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.cxx b/GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.cxx index 30a8fc193774b..03f123e97fb78 100644 --- a/GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.cxx +++ b/GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.cxx @@ -110,6 +110,45 @@ int32_t GPUReconstructionOCLBackend::InitDevice_Runtime() } bool found = false; + char platform_profile[256] = {}, platform_version[256] = {}, platform_name[256] = {}, platform_vendor[256] = {}; + auto queryPlatforms = [&platform_profile, &platform_version, &platform_name, &platform_vendor](auto platform) { + clGetPlatformInfo(platform, CL_PLATFORM_PROFILE, sizeof(platform_profile), platform_profile, nullptr); + clGetPlatformInfo(platform, CL_PLATFORM_VERSION, sizeof(platform_version), platform_version, nullptr); + clGetPlatformInfo(platform, CL_PLATFORM_NAME, sizeof(platform_name), platform_name, nullptr); + clGetPlatformInfo(platform, CL_PLATFORM_VENDOR, sizeof(platform_vendor), platform_vendor, nullptr); + }; + auto checkPlatform = [&](auto platform) { + cl_uint tmp; + if (clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 0, nullptr, &tmp) != CL_SUCCESS || tmp == 0) { + return false; + } + + queryPlatforms(platform); + float ver1 = 0; + sscanf(platform_version, "OpenCL %f", &ver1); + if (ver1 >= 2.2f) { + if (mProcessingSettings.debugLevel >= 2) { + GPUInfo("OpenCL 2.2 capable platform found"); + } + return true; + } + + if (strcmp(platform_vendor, "Advanced Micro Devices, Inc.") == 0 && ver1 >= 2.0f) { + float ver2 = 0; + const char* pos = strchr(platform_version, '('); + if (pos) { + sscanf(pos, "(%f)", &ver2); + } + if ((ver1 >= 2.f && ver2 >= 2000.f) || ver1 >= 2.1f) { + if (mProcessingSettings.debugLevel >= 2) { + GPUInfo("AMD ROCm OpenCL Platform found"); + } + return true; + } + } + return false; + }; + if (mProcessingSettings.platformNum >= 0) { if (mProcessingSettings.platformNum >= (int32_t)num_platforms) { quit("Invalid platform specified"); @@ -117,22 +156,14 @@ int32_t GPUReconstructionOCLBackend::InitDevice_Runtime() mInternals->platform = mInternals->platforms[mProcessingSettings.platformNum]; found = true; if (mProcessingSettings.debugLevel >= 2) { - char platform_profile[256] = {}, platform_version[256] = {}, platform_name[256] = {}, platform_vendor[256] = {}; - clGetPlatformInfo(mInternals->platform, CL_PLATFORM_PROFILE, sizeof(platform_profile), platform_profile, nullptr); - clGetPlatformInfo(mInternals->platform, CL_PLATFORM_VERSION, sizeof(platform_version), platform_version, nullptr); - clGetPlatformInfo(mInternals->platform, CL_PLATFORM_NAME, sizeof(platform_name), platform_name, nullptr); - clGetPlatformInfo(mInternals->platform, CL_PLATFORM_VENDOR, sizeof(platform_vendor), platform_vendor, nullptr); + queryPlatforms(mInternals->platform); GPUInfo("Selected Platform %d: (%s %s) %s %s", mProcessingSettings.platformNum, platform_profile, platform_version, platform_vendor, platform_name); } } else { for (uint32_t i_platform = 0; i_platform < num_platforms; i_platform++) { - char platform_profile[256] = {}, platform_version[256] = {}, platform_name[256] = {}, platform_vendor[256] = {}; - clGetPlatformInfo(mInternals->platforms[i_platform], CL_PLATFORM_PROFILE, sizeof(platform_profile), platform_profile, nullptr); - clGetPlatformInfo(mInternals->platforms[i_platform], CL_PLATFORM_VERSION, sizeof(platform_version), platform_version, nullptr); - clGetPlatformInfo(mInternals->platforms[i_platform], CL_PLATFORM_NAME, sizeof(platform_name), platform_name, nullptr); - clGetPlatformInfo(mInternals->platforms[i_platform], CL_PLATFORM_VENDOR, sizeof(platform_vendor), platform_vendor, nullptr); + queryPlatforms(mInternals->platforms[i_platform]); const char* platformUsageInfo = ""; - if (!found && CheckPlatform(i_platform)) { + if (!found && checkPlatform(mInternals->platforms[i_platform])) { found = true; mInternals->platform = mInternals->platforms[i_platform]; if (mProcessingSettings.debugLevel >= 2) { @@ -149,14 +180,14 @@ int32_t GPUReconstructionOCLBackend::InitDevice_Runtime() quit("Did not find compatible OpenCL Platform"); } - cl_uint count, bestDevice = (cl_uint)-1; - if (GPUFailedMsgI(clGetDeviceIDs(mInternals->platform, CL_DEVICE_TYPE_ALL, 0, nullptr, &count))) { + cl_uint deviceCount, bestDevice = (cl_uint)-1; + if (GPUFailedMsgI(clGetDeviceIDs(mInternals->platform, CL_DEVICE_TYPE_ALL, 0, nullptr, &deviceCount))) { quit("Error getting OPENCL Device Count"); } // Query devices - mInternals->devices.reset(new cl_device_id[count]); - if (GPUFailedMsgI(clGetDeviceIDs(mInternals->platform, CL_DEVICE_TYPE_ALL, count, mInternals->devices.get(), nullptr))) { + mInternals->devices.reset(new cl_device_id[deviceCount]); + if (GPUFailedMsgI(clGetDeviceIDs(mInternals->platform, CL_DEVICE_TYPE_ALL, deviceCount, mInternals->devices.get(), nullptr))) { quit("Error getting OpenCL devices"); } @@ -167,8 +198,8 @@ int32_t GPUReconstructionOCLBackend::InitDevice_Runtime() if (mProcessingSettings.debugLevel >= 2) { GPUInfo("Available OPENCL devices:"); } - std::vector devicesOK(count, false); - for (uint32_t i = 0; i < count; i++) { + std::vector devicesOK(deviceCount, false); + for (uint32_t i = 0; i < deviceCount; i++) { if (mProcessingSettings.debugLevel >= 3) { GPUInfo("Examining device %d", i); } @@ -215,11 +246,11 @@ int32_t GPUReconstructionOCLBackend::InitDevice_Runtime() } } if (bestDevice == (cl_uint)-1) { - quit("No %sOPENCL Device available, aborting OPENCL Initialisation", count ? "appropriate " : ""); + quit("No %sOPENCL Device available, aborting OPENCL Initialisation", deviceCount ? "appropriate " : ""); } if (mProcessingSettings.deviceNum > -1) { - if (mProcessingSettings.deviceNum >= (signed)count) { + if (mProcessingSettings.deviceNum >= (signed)deviceCount) { quit("Requested device ID %d does not exist", mProcessingSettings.deviceNum); } else if (!devicesOK[mProcessingSettings.deviceNum]) { quit("Unsupported device requested (%d)", mProcessingSettings.deviceNum); @@ -269,7 +300,7 @@ int32_t GPUReconstructionOCLBackend::InitDevice_Runtime() mWarpSize = 32; mMaxBackendThreads = std::max(mMaxBackendThreads, maxWorkGroup * mBlockCount); - mInternals->context = clCreateContext(nullptr, ContextForAllPlatforms() ? count : 1, ContextForAllPlatforms() ? mInternals->devices.get() : &mInternals->device, nullptr, nullptr, &ocl_error); + mInternals->context = clCreateContext(nullptr, ContextForAllPlatforms() ? deviceCount : 1, ContextForAllPlatforms() ? mInternals->devices.get() : &mInternals->device, nullptr, nullptr, &ocl_error); if (GPUFailedMsgI(ocl_error)) { quit("Could not create OPENCL Device Context!"); } @@ -608,33 +639,3 @@ int32_t GPUReconstructionOCLBackend::GetOCLPrograms() return 0; } - -bool GPUReconstructionOCLBackend::CheckPlatform(uint32_t i) -{ - char platform_version[64] = {}, platform_vendor[64] = {}; - clGetPlatformInfo(mInternals->platforms[i], CL_PLATFORM_VERSION, sizeof(platform_version), platform_version, nullptr); - clGetPlatformInfo(mInternals->platforms[i], CL_PLATFORM_VENDOR, sizeof(platform_vendor), platform_vendor, nullptr); - float ver1 = 0; - sscanf(platform_version, "OpenCL %f", &ver1); - if (ver1 >= 2.2f) { - if (mProcessingSettings.debugLevel >= 2) { - GPUInfo("OpenCL 2.2 capable platform found"); - } - return true; - } - - if (strcmp(platform_vendor, "Advanced Micro Devices, Inc.") == 0 && ver1 >= 2.0f) { - float ver2 = 0; - const char* pos = strchr(platform_version, '('); - if (pos) { - sscanf(pos, "(%f)", &ver2); - } - if ((ver1 >= 2.f && ver2 >= 2000.f) || ver1 >= 2.1f) { - if (mProcessingSettings.debugLevel >= 2) { - GPUInfo("AMD ROCm OpenCL Platform found"); - } - return true; - } - } - return false; -} diff --git a/GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.h b/GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.h index 245e9674801f6..d052ba53dfac5 100644 --- a/GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.h +++ b/GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.h @@ -74,7 +74,6 @@ class GPUReconstructionOCLBackend : public GPUReconstructionDeviceBase S& getKernelObject(); int32_t GetOCLPrograms(); - bool CheckPlatform(uint32_t i); }; using GPUReconstructionOCL = GPUReconstructionKernels; From 887e18066a08e833176888d49d871bce1b514172 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Mon, 24 Feb 2025 23:01:07 +0100 Subject: [PATCH 2/3] GPU OpenCL: Improve device detection and do not consider platforms/devices, which do not support SPIR-V --- GPU/GPUTracking/Base/opencl/CMakeLists.txt | 2 +- .../Base/opencl/GPUReconstructionOCL.cxx | 382 +++++++++--------- .../Base/opencl/GPUReconstructionOCL.h | 3 +- .../opencl/GPUReconstructionOCLInternals.h | 2 - GPU/GPUTracking/Definitions/GPUSettingsList.h | 3 +- 5 files changed, 187 insertions(+), 205 deletions(-) diff --git a/GPU/GPUTracking/Base/opencl/CMakeLists.txt b/GPU/GPUTracking/Base/opencl/CMakeLists.txt index 1aa3739b0b44a..2a361356283a8 100644 --- a/GPU/GPUTracking/Base/opencl/CMakeLists.txt +++ b/GPU/GPUTracking/Base/opencl/CMakeLists.txt @@ -106,4 +106,4 @@ endif() if(OPENCL_ENABLED_SPIRV) target_compile_definitions(${targetName} PRIVATE OPENCL_ENABLED_SPIRV) endif() -target_compile_definitions(${targetName} PRIVATE OCL_FLAGS=$) +target_compile_definitions(${targetName} PRIVATE GPUCA_OCL_BUILD_FLAGS=$) diff --git a/GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.cxx b/GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.cxx index 03f123e97fb78..6b918fe501330 100644 --- a/GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.cxx +++ b/GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.cxx @@ -26,10 +26,10 @@ using namespace o2::gpu; #include #include -#define quit(...) \ - { \ - GPUError(__VA_ARGS__); \ - return (1); \ +#define GPUErrorReturn(...) \ + { \ + GPUError(__VA_ARGS__); \ + return (1); \ } #define GPUCA_KRNL(x_class, x_attributes, ...) GPUCA_KRNL_PROP(x_class, x_attributes) @@ -94,215 +94,205 @@ int32_t GPUReconstructionOCLBackend::InitDevice_Runtime() cl_int ocl_error; cl_uint num_platforms; if (GPUFailedMsgI(clGetPlatformIDs(0, nullptr, &num_platforms))) { - quit("Error getting OpenCL Platform Count"); + GPUErrorReturn("Error getting OpenCL Platform Count"); } if (num_platforms == 0) { - quit("No OpenCL Platform found"); + GPUErrorReturn("No OpenCL Platform found"); } if (mProcessingSettings.debugLevel >= 2) { GPUInfo("%d OpenCL Platforms found", num_platforms); } - // Query platforms - mInternals->platforms.reset(new cl_platform_id[num_platforms]); - if (GPUFailedMsgI(clGetPlatformIDs(num_platforms, mInternals->platforms.get(), nullptr))) { - quit("Error getting OpenCL Platforms"); + // Query platforms and devices + std::unique_ptr platforms; + platforms.reset(new cl_platform_id[num_platforms]); + if (GPUFailedMsgI(clGetPlatformIDs(num_platforms, platforms.get(), nullptr))) { + GPUErrorReturn("Error getting OpenCL Platforms"); } - bool found = false; - char platform_profile[256] = {}, platform_version[256] = {}, platform_name[256] = {}, platform_vendor[256] = {}; - auto queryPlatforms = [&platform_profile, &platform_version, &platform_name, &platform_vendor](auto platform) { - clGetPlatformInfo(platform, CL_PLATFORM_PROFILE, sizeof(platform_profile), platform_profile, nullptr); - clGetPlatformInfo(platform, CL_PLATFORM_VERSION, sizeof(platform_version), platform_version, nullptr); - clGetPlatformInfo(platform, CL_PLATFORM_NAME, sizeof(platform_name), platform_name, nullptr); - clGetPlatformInfo(platform, CL_PLATFORM_VENDOR, sizeof(platform_vendor), platform_vendor, nullptr); + auto query = [&](auto func, auto obj, auto var) { + size_t size; + func(obj, var, 0, nullptr, &size); + std::string retVal(size - 1, ' '); + func(obj, var, size, retVal.data(), nullptr); + return retVal; }; - auto checkPlatform = [&](auto platform) { - cl_uint tmp; - if (clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 0, nullptr, &tmp) != CL_SUCCESS || tmp == 0) { - return false; - } - queryPlatforms(platform); - float ver1 = 0; - sscanf(platform_version, "OpenCL %f", &ver1); - if (ver1 >= 2.2f) { - if (mProcessingSettings.debugLevel >= 2) { - GPUInfo("OpenCL 2.2 capable platform found"); - } - return true; - } + std::string platform_profile, platform_version, platform_name, platform_vendor; + float platform_version_f; + auto queryPlatform = [&](auto platform) { + platform_profile = query(clGetPlatformInfo, platform, CL_PLATFORM_PROFILE); + platform_version = query(clGetPlatformInfo, platform, CL_PLATFORM_VERSION); + platform_name = query(clGetPlatformInfo, platform, CL_PLATFORM_NAME); + platform_vendor = query(clGetPlatformInfo, platform, CL_PLATFORM_VENDOR); + sscanf(platform_version.c_str(), "OpenCL %f", &platform_version_f); + }; - if (strcmp(platform_vendor, "Advanced Micro Devices, Inc.") == 0 && ver1 >= 2.0f) { - float ver2 = 0; - const char* pos = strchr(platform_version, '('); - if (pos) { - sscanf(pos, "(%f)", &ver2); - } - if ((ver1 >= 2.f && ver2 >= 2000.f) || ver1 >= 2.1f) { - if (mProcessingSettings.debugLevel >= 2) { - GPUInfo("AMD ROCm OpenCL Platform found"); - } - return true; - } - } - return false; + std::vector devices; + std::string device_vendor, device_name, device_il_version; + cl_device_type device_type; + cl_uint device_freq, device_shaders, device_nbits; + cl_bool device_endian; + auto queryDevice = [&](auto device) { + platform_name = query(clGetDeviceInfo, device, CL_DEVICE_NAME); + device_vendor = query(clGetDeviceInfo, device, CL_DEVICE_VENDOR); + device_il_version = query(clGetDeviceInfo, device, CL_DEVICE_IL_VERSION); + clGetDeviceInfo(device, CL_DEVICE_TYPE, sizeof(device_type), &device_type, nullptr); + clGetDeviceInfo(device, CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof(device_freq), &device_freq, nullptr); + clGetDeviceInfo(device, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(device_shaders), &device_shaders, nullptr); + clGetDeviceInfo(device, CL_DEVICE_ADDRESS_BITS, sizeof(device_nbits), &device_nbits, nullptr); + clGetDeviceInfo(device, CL_DEVICE_ENDIAN_LITTLE, sizeof(device_endian), &device_endian, nullptr); }; - if (mProcessingSettings.platformNum >= 0) { - if (mProcessingSettings.platformNum >= (int32_t)num_platforms) { - quit("Invalid platform specified"); - } - mInternals->platform = mInternals->platforms[mProcessingSettings.platformNum]; - found = true; - if (mProcessingSettings.debugLevel >= 2) { - queryPlatforms(mInternals->platform); - GPUInfo("Selected Platform %d: (%s %s) %s %s", mProcessingSettings.platformNum, platform_profile, platform_version, platform_vendor, platform_name); + cl_uint deviceCount, bestDevice = (cl_uint)-1, bestPlatform = (cl_uint)-1; + for (uint32_t iPlatform = 0; iPlatform < num_platforms; iPlatform++) { + if (mProcessingSettings.oclPlatformNum >= 0) { + if (mProcessingSettings.oclPlatformNum >= (int32_t)num_platforms) { + GPUErrorReturn("Invalid platform specified"); + } + iPlatform = mProcessingSettings.oclPlatformNum; } - } else { - for (uint32_t i_platform = 0; i_platform < num_platforms; i_platform++) { - queryPlatforms(mInternals->platforms[i_platform]); - const char* platformUsageInfo = ""; - if (!found && checkPlatform(mInternals->platforms[i_platform])) { - found = true; - mInternals->platform = mInternals->platforms[i_platform]; - if (mProcessingSettings.debugLevel >= 2) { - platformUsageInfo = " !!! Using this platform !!!"; - } + std::string platformUsageInfo; + bool platformCompatible = false; + queryPlatform(platforms[iPlatform]); + if (clGetDeviceIDs(platforms[iPlatform], CL_DEVICE_TYPE_ALL, 0, nullptr, &deviceCount) != CL_SUCCESS) { + if (mProcessingSettings.oclPlatformNum >= 0) { + GPUErrorReturn("No device in requested platform or error obtaining device count"); } - if (mProcessingSettings.debugLevel >= 2) { - GPUInfo("Available Platform %d: (%s %s) %s %s%s", i_platform, platform_profile, platform_version, platform_vendor, platform_name, platformUsageInfo); + platformUsageInfo += " - no devices"; + } else { + if (platform_version_f >= 2.1f) { + platformUsageInfo += " - OpenCL 2.2 capable"; + platformCompatible = true; } } - } - - if (found == false) { - quit("Did not find compatible OpenCL Platform"); - } - - cl_uint deviceCount, bestDevice = (cl_uint)-1; - if (GPUFailedMsgI(clGetDeviceIDs(mInternals->platform, CL_DEVICE_TYPE_ALL, 0, nullptr, &deviceCount))) { - quit("Error getting OPENCL Device Count"); - } - // Query devices - mInternals->devices.reset(new cl_device_id[deviceCount]); - if (GPUFailedMsgI(clGetDeviceIDs(mInternals->platform, CL_DEVICE_TYPE_ALL, deviceCount, mInternals->devices.get(), nullptr))) { - quit("Error getting OpenCL devices"); - } + if (mProcessingSettings.oclPlatformNum >= 0 || mProcessingSettings.debugLevel >= 2) { + GPUInfo("%s Platform %d: (%s %s) %s %s (Compatible: %s)%s", mProcessingSettings.oclPlatformNum >= 0 ? "Enforced" : "Available", iPlatform, platform_profile.c_str(), platform_version.c_str(), platform_vendor.c_str(), platform_name.c_str(), platformCompatible ? "yes" : "no", mProcessingSettings.debugLevel >= 2 ? platformUsageInfo.c_str() : ""); + } - char device_vendor[64], device_name[64]; - cl_device_type device_type; - cl_uint freq, shaders; + if (platformCompatible || mProcessingSettings.oclPlatformNum >= 0 || (mProcessingSettings.oclPlatformNum == -2 && deviceCount)) { + if (deviceCount > devices.size()) { + devices.resize(deviceCount); + } + if (clGetDeviceIDs(platforms[iPlatform], CL_DEVICE_TYPE_ALL, deviceCount, devices.data(), nullptr) != CL_SUCCESS) { + if (mProcessingSettings.oclPlatformNum >= 0) { + GPUErrorReturn("Error getting OpenCL devices"); + } + continue; + } - if (mProcessingSettings.debugLevel >= 2) { - GPUInfo("Available OPENCL devices:"); - } - std::vector devicesOK(deviceCount, false); - for (uint32_t i = 0; i < deviceCount; i++) { - if (mProcessingSettings.debugLevel >= 3) { - GPUInfo("Examining device %d", i); - } - cl_uint nbits; - cl_bool endian; - - clGetDeviceInfo(mInternals->devices[i], CL_DEVICE_NAME, 64, device_name, nullptr); - clGetDeviceInfo(mInternals->devices[i], CL_DEVICE_VENDOR, 64, device_vendor, nullptr); - clGetDeviceInfo(mInternals->devices[i], CL_DEVICE_TYPE, sizeof(cl_device_type), &device_type, nullptr); - clGetDeviceInfo(mInternals->devices[i], CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof(freq), &freq, nullptr); - clGetDeviceInfo(mInternals->devices[i], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(shaders), &shaders, nullptr); - clGetDeviceInfo(mInternals->devices[i], CL_DEVICE_ADDRESS_BITS, sizeof(nbits), &nbits, nullptr); - clGetDeviceInfo(mInternals->devices[i], CL_DEVICE_ENDIAN_LITTLE, sizeof(endian), &endian, nullptr); - int32_t deviceOK = true; - const char* deviceFailure = ""; - if (mProcessingSettings.gpuDeviceOnly && ((device_type & CL_DEVICE_TYPE_CPU) || !(device_type & CL_DEVICE_TYPE_GPU))) { - deviceOK = false; - deviceFailure = "No GPU device"; - } - if (nbits / 8 != sizeof(void*)) { - deviceOK = false; - deviceFailure = "No 64 bit device"; - } - if (!endian) { - deviceOK = false; - deviceFailure = "No Little Endian Mode"; - } + for (uint32_t i = 0; i < deviceCount; i++) { + if (mProcessingSettings.deviceNum >= 0) { + if (mProcessingSettings.deviceNum >= (signed)deviceCount) { + GPUErrorReturn("Requested device ID %d does not exist", mProcessingSettings.deviceNum); + } + i = mProcessingSettings.deviceNum; + } + bool deviceOK = true; + queryDevice(devices[i]); + std::string deviceFailure; + if (mProcessingSettings.gpuDeviceOnly && ((device_type & CL_DEVICE_TYPE_CPU) || !(device_type & CL_DEVICE_TYPE_GPU))) { + deviceOK = false; + deviceFailure += " - No GPU device"; + } + if (device_nbits / 8 != sizeof(void*)) { + deviceOK = false; + deviceFailure += " - No 64 bit device"; + } + if (!device_endian) { + deviceOK = false; + deviceFailure += " - No Little Endian Mode"; + } + if (!GetProcessingSettings().oclCompileFromSources) { + size_t pos = 0; + while ((pos = device_il_version.find("SPIR-V", pos)) != std::string::npos) { + float spirvVersion; + sscanf(device_il_version.c_str() + pos, "SPIR-V_%f", &spirvVersion); + if (spirvVersion >= 1.2) { + break; + } + pos += strlen("SPIR-V_0.0"); + } + if (pos == std::string::npos) { + deviceOK = false; + deviceFailure += " - No SPIR-V 1.6 (" + device_il_version + ")"; + } + } - double bestDeviceSpeed = -1, deviceSpeed = (double)freq * (double)shaders; - if (mProcessingSettings.debugLevel >= 2) { - GPUImportant("Device %s%2d: %s %s (Frequency %d, Shaders %d, %d bit) (Speed Value: %ld)%s %s", deviceOK ? " " : "[", i, device_vendor, device_name, (int32_t)freq, (int32_t)shaders, (int32_t)nbits, (int64_t)deviceSpeed, deviceOK ? " " : " ]", deviceOK ? "" : deviceFailure); - } - if (!deviceOK) { - continue; - } - devicesOK[i] = true; - if (deviceSpeed > bestDeviceSpeed) { - bestDevice = i; - bestDeviceSpeed = deviceSpeed; - } else { - if (mProcessingSettings.debugLevel >= 2) { - GPUInfo("Skipping: Speed %f < %f", deviceSpeed, bestDeviceSpeed); + double bestDeviceSpeed = -1, deviceSpeed = (double)device_freq * (double)device_shaders; + if (mProcessingSettings.debugLevel >= 2) { + GPUInfo(" Device %s%2d: %s %s (Frequency %d, Shaders %d, %d bit) (Speed Value: %ld)%s %s", deviceOK ? " " : "[", i, device_vendor.c_str(), device_name.c_str(), (int32_t)device_freq, (int32_t)device_shaders, (int32_t)device_nbits, (int64_t)deviceSpeed, deviceOK ? " " : " ]", deviceOK ? "" : deviceFailure.c_str()); + } + if (!deviceOK) { + if (mProcessingSettings.deviceNum >= 0) { + GPUInfo("Unsupported device requested on platform %d: (%d)", iPlatform, mProcessingSettings.deviceNum); + break; + } + continue; + } + if (deviceSpeed > bestDeviceSpeed) { + bestDevice = i; + bestPlatform = iPlatform; + bestDeviceSpeed = deviceSpeed; + mOclVersion = platform_version_f; + } + if (mProcessingSettings.deviceNum >= 0) { + break; + } } } - } - if (bestDevice == (cl_uint)-1) { - quit("No %sOPENCL Device available, aborting OPENCL Initialisation", deviceCount ? "appropriate " : ""); - } - - if (mProcessingSettings.deviceNum > -1) { - if (mProcessingSettings.deviceNum >= (signed)deviceCount) { - quit("Requested device ID %d does not exist", mProcessingSettings.deviceNum); - } else if (!devicesOK[mProcessingSettings.deviceNum]) { - quit("Unsupported device requested (%d)", mProcessingSettings.deviceNum); - } else { - bestDevice = mProcessingSettings.deviceNum; + if (mProcessingSettings.oclPlatformNum >= 0) { + break; } } - mInternals->device = mInternals->devices[bestDevice]; - - cl_ulong constantBuffer, globalMem, localMem; - char deviceVersion[64]; - size_t maxWorkGroup, maxWorkItems[3]; - clGetDeviceInfo(mInternals->device, CL_DEVICE_NAME, 64, device_name, nullptr); - clGetDeviceInfo(mInternals->device, CL_DEVICE_VENDOR, 64, device_vendor, nullptr); - clGetDeviceInfo(mInternals->device, CL_DEVICE_TYPE, sizeof(cl_device_type), &device_type, nullptr); - clGetDeviceInfo(mInternals->device, CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof(freq), &freq, nullptr); - clGetDeviceInfo(mInternals->device, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(shaders), &shaders, nullptr); - clGetDeviceInfo(mInternals->device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(globalMem), &globalMem, nullptr); - clGetDeviceInfo(mInternals->device, CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, sizeof(constantBuffer), &constantBuffer, nullptr); - clGetDeviceInfo(mInternals->device, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(localMem), &localMem, nullptr); - clGetDeviceInfo(mInternals->device, CL_DEVICE_VERSION, sizeof(deviceVersion) - 1, deviceVersion, nullptr); - clGetDeviceInfo(mInternals->device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(maxWorkGroup), &maxWorkGroup, nullptr); - clGetDeviceInfo(mInternals->device, CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(maxWorkItems), maxWorkItems, nullptr); + + if (bestDevice == (cl_uint)-1) { + GPUErrorReturn("Did not find compatible OpenCL Platform / Device, aborting OPENCL Initialisation"); + } + mInternals->platform = platforms[bestPlatform]; + GPUFailedMsg(clGetDeviceIDs(mInternals->platform, CL_DEVICE_TYPE_ALL, devices.size(), devices.data(), nullptr)); + mInternals->device = devices[bestDevice]; + queryDevice(mInternals->device); + + cl_ulong deviceConstantBuffer, deviceGlobalMem, deviceLocalMem; + std::string deviceVersion; + size_t deviceMaxWorkGroup, deviceMaxWorkItems[3]; + clGetDeviceInfo(mInternals->device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(deviceGlobalMem), &deviceGlobalMem, nullptr); + clGetDeviceInfo(mInternals->device, CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, sizeof(deviceConstantBuffer), &deviceConstantBuffer, nullptr); + clGetDeviceInfo(mInternals->device, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(deviceLocalMem), &deviceLocalMem, nullptr); + clGetDeviceInfo(mInternals->device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(deviceMaxWorkGroup), &deviceMaxWorkGroup, nullptr); + clGetDeviceInfo(mInternals->device, CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(deviceMaxWorkItems), deviceMaxWorkItems, nullptr); + deviceVersion = query(clGetDeviceInfo, mInternals->device, CL_DEVICE_VERSION); int versionMajor, versionMinor; - sscanf(deviceVersion, "OpenCL %d.%d", &versionMajor, &versionMinor); + sscanf(deviceVersion.c_str(), "OpenCL %d.%d", &versionMajor, &versionMinor); if (mProcessingSettings.debugLevel >= 2) { - GPUInfo("Using OpenCL device %d: %s %s with properties:", bestDevice, device_vendor, device_name); + GPUInfo("Using OpenCL platform %d / device %d: %s %s with properties:", bestPlatform, bestDevice, device_vendor.c_str(), device_name.c_str()); GPUInfo("\tVersion = %s", deviceVersion); - GPUInfo("\tFrequency = %d", (int32_t)freq); - GPUInfo("\tShaders = %d", (int32_t)shaders); - GPUInfo("\tGLobalMemory = %ld", (int64_t)globalMem); - GPUInfo("\tContantMemoryBuffer = %ld", (int64_t)constantBuffer); - GPUInfo("\tLocalMemory = %ld", (int64_t)localMem); - GPUInfo("\tmaxThreadsPerBlock = %ld", (int64_t)maxWorkGroup); - GPUInfo("\tmaxThreadsDim = %ld %ld %ld", (int64_t)maxWorkItems[0], (int64_t)maxWorkItems[1], (int64_t)maxWorkItems[2]); + GPUInfo("\tFrequency = %d", (int32_t)device_freq); + GPUInfo("\tShaders = %d", (int32_t)device_shaders); + GPUInfo("\tGLobalMemory = %ld", (int64_t)deviceGlobalMem); + GPUInfo("\tContantMemoryBuffer = %ld", (int64_t)deviceConstantBuffer); + GPUInfo("\tLocalMemory = %ld", (int64_t)deviceLocalMem); + GPUInfo("\tmaxThreadsPerBlock = %ld", (int64_t)deviceMaxWorkGroup); + GPUInfo("\tmaxThreadsDim = %ld %ld %ld", (int64_t)deviceMaxWorkItems[0], (int64_t)deviceMaxWorkItems[1], (int64_t)deviceMaxWorkItems[2]); GPUInfo(" "); } #ifndef GPUCA_NO_CONSTANT_MEMORY - if (gGPUConstantMemBufferSize > constantBuffer) { - quit("Insufficient constant memory available on GPU %d < %d!", (int32_t)constantBuffer, (int32_t)gGPUConstantMemBufferSize); + if (gGPUConstantMemBufferSize > deviceConstantBuffer) { + GPUErrorReturn("Insufficient constant memory available on GPU %d < %d!", (int32_t)deviceConstantBuffer, (int32_t)gGPUConstantMemBufferSize); } #endif - mDeviceName = device_name; + mDeviceName = device_name.c_str(); mDeviceName += " (OpenCL)"; - mBlockCount = shaders; + mBlockCount = device_shaders; mWarpSize = 32; - mMaxBackendThreads = std::max(mMaxBackendThreads, maxWorkGroup * mBlockCount); + mMaxBackendThreads = std::max(mMaxBackendThreads, deviceMaxWorkGroup * mBlockCount); - mInternals->context = clCreateContext(nullptr, ContextForAllPlatforms() ? deviceCount : 1, ContextForAllPlatforms() ? mInternals->devices.get() : &mInternals->device, nullptr, nullptr, &ocl_error); + mInternals->context = clCreateContext(nullptr, 1, &mInternals->device, nullptr, nullptr, &ocl_error); if (GPUFailedMsgI(ocl_error)) { - quit("Could not create OPENCL Device Context!"); + GPUErrorReturn("Could not create OPENCL Device Context!"); } if (GetOCLPrograms()) { @@ -316,14 +306,14 @@ int32_t GPUReconstructionOCLBackend::InitDevice_Runtime() mInternals->mem_gpu = clCreateBuffer(mInternals->context, CL_MEM_READ_WRITE, mDeviceMemorySize, nullptr, &ocl_error); if (GPUFailedMsgI(ocl_error)) { clReleaseContext(mInternals->context); - quit("OPENCL Memory Allocation Error"); + GPUErrorReturn("OPENCL Memory Allocation Error"); } mInternals->mem_constant = clCreateBuffer(mInternals->context, CL_MEM_READ_ONLY, gGPUConstantMemBufferSize, nullptr, &ocl_error); if (GPUFailedMsgI(ocl_error)) { clReleaseMemObject(mInternals->mem_gpu); clReleaseContext(mInternals->context); - quit("OPENCL Constant Memory Allocation Error"); + GPUErrorReturn("OPENCL Constant Memory Allocation Error"); } if (device_type & CL_DEVICE_TYPE_CPU) { @@ -349,36 +339,36 @@ int32_t GPUReconstructionOCLBackend::InitDevice_Runtime() mInternals->command_queue[i] = clCreateCommandQueue(mInternals->context, mInternals->device, 0, &ocl_error); #endif if (GPUFailedMsgI(ocl_error)) { - quit("Error creating OpenCL command queue"); + GPUErrorReturn("Error creating OpenCL command queue"); } } if (GPUFailedMsgI(clEnqueueMigrateMemObjects(mInternals->command_queue[0], 1, &mInternals->mem_gpu, 0, 0, nullptr, nullptr))) { - quit("Error migrating buffer"); + GPUErrorReturn("Error migrating buffer"); } if (GPUFailedMsgI(clEnqueueMigrateMemObjects(mInternals->command_queue[0], 1, &mInternals->mem_constant, 0, 0, nullptr, nullptr))) { - quit("Error migrating buffer"); + GPUErrorReturn("Error migrating buffer"); } mInternals->mem_host = clCreateBuffer(mInternals->context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, mHostMemorySize, nullptr, &ocl_error); if (GPUFailedMsgI(ocl_error)) { - quit("Error allocating pinned host memory"); + GPUErrorReturn("Error allocating pinned host memory"); } const char* krnlGetPtr = "__kernel void krnlGetPtr(__global char* gpu_mem, __global char* constant_mem, __global size_t* host_mem) {if (get_global_id(0) == 0) {host_mem[0] = (size_t) gpu_mem; host_mem[1] = (size_t) constant_mem;}}"; cl_program program = clCreateProgramWithSource(mInternals->context, 1, (const char**)&krnlGetPtr, nullptr, &ocl_error); if (GPUFailedMsgI(ocl_error)) { - quit("Error creating program object"); + GPUErrorReturn("Error creating program object"); } ocl_error = clBuildProgram(program, 1, &mInternals->device, "", nullptr, nullptr); if (GPUFailedMsgI(ocl_error)) { char build_log[16384]; clGetProgramBuildInfo(program, mInternals->device, CL_PROGRAM_BUILD_LOG, 16384, build_log, nullptr); GPUImportant("Build Log:\n\n%s\n\n", build_log); - quit("Error compiling program"); + GPUErrorReturn("Error compiling program"); } cl_kernel kernel = clCreateKernel(program, "krnlGetPtr", &ocl_error); if (GPUFailedMsgI(ocl_error)) { - quit("Error creating kernel"); + GPUErrorReturn("Error creating kernel"); } if (GPUFailedMsgI(OCLsetKernelParameters(kernel, mInternals->mem_gpu, mInternals->mem_constant, mInternals->mem_host)) || @@ -386,7 +376,7 @@ int32_t GPUReconstructionOCLBackend::InitDevice_Runtime() GPUFailedMsgI(clFinish(mInternals->command_queue[0])) || GPUFailedMsgI(clReleaseKernel(kernel)) || GPUFailedMsgI(clReleaseProgram(program))) { - quit("Error obtaining device memory ptr"); + GPUErrorReturn("Error obtaining device memory ptr"); } if (mProcessingSettings.debugLevel >= 2) { @@ -394,7 +384,7 @@ int32_t GPUReconstructionOCLBackend::InitDevice_Runtime() } mHostMemoryBase = clEnqueueMapBuffer(mInternals->command_queue[0], mInternals->mem_host, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, mHostMemorySize, 0, nullptr, nullptr, &ocl_error); if (GPUFailedMsgI(ocl_error)) { - quit("Error allocating Page Locked Host Memory"); + GPUErrorReturn("Error allocating Page Locked Host Memory"); } mDeviceMemoryBase = ((void**)mHostMemoryBase)[0]; @@ -405,8 +395,7 @@ int32_t GPUReconstructionOCLBackend::InitDevice_Runtime() memset(mHostMemoryBase, 0xDD, mHostMemorySize); } - GPUInfo("OPENCL Initialisation successfull (%d: %s %s (Frequency %d, Shaders %d), %ld / %ld bytes host / global memory, Stack frame %d, Constant memory %ld)", bestDevice, device_vendor, device_name, (int32_t)freq, (int32_t)shaders, (int64_t)mDeviceMemorySize, - (int64_t)mHostMemorySize, -1, (int64_t)gGPUConstantMemBufferSize); + GPUInfo("OPENCL Initialisation successfull (%d: %s %s (Frequency %d, Shaders %d), %ld / %ld bytes host / global memory, Stack frame %d, Constant memory %ld)", bestDevice, device_vendor, device_name, (int32_t)device_freq, (int32_t)device_shaders, (int64_t)mDeviceMemorySize, (int64_t)mHostMemorySize, -1, (int64_t)gGPUConstantMemBufferSize); } else { GPUReconstructionOCL* master = dynamic_cast(mMaster); mBlockCount = master->mBlockCount; @@ -510,7 +499,7 @@ int32_t GPUReconstructionOCLBackend::DoStuckProtection(int32_t stream, deviceEve } if (tmp != CL_COMPLETE) { mGPUStuck = 1; - quit("GPU Stuck, future processing in this component is disabled, skipping event (GPU Event State %d)", (int32_t)tmp); + GPUErrorReturn("GPU Stuck, future processing in this component is disabled, skipping event (GPU Event State %d)", (int32_t)tmp); } } else { clFinish(mInternals->command_queue[stream]); @@ -581,24 +570,19 @@ S& GPUReconstructionOCLBackend::getKernelObject() int32_t GPUReconstructionOCLBackend::GetOCLPrograms() { - char platform_version[256] = {}; - GPUFailedMsg(clGetPlatformInfo(mInternals->platform, CL_PLATFORM_VERSION, sizeof(platform_version), platform_version, nullptr)); - float ver = 0; - sscanf(platform_version, "OpenCL %f", &ver); - cl_int ocl_error; - const char* ocl_flags = GPUCA_M_STR(OCL_FLAGS); + const char* oclBuildFlags = GetProcessingSettings().oclOverrideSourceBuildFlags != "" ? GetProcessingSettings().oclOverrideSourceBuildFlags.c_str() : GPUCA_M_STR(GPUCA_OCL_BUILD_FLAGS); #ifdef OPENCL_ENABLED_SPIRV // clang-format off - if (ver >= 2.2f && !GetProcessingSettings().oclCompileFromSources) { - GPUInfo("Reading OpenCL program from SPIR-V IL (Platform version %4.2f)", ver); + if (mOclVersion >= 2.1f && !GetProcessingSettings().oclCompileFromSources) { + GPUInfo("Reading OpenCL program from SPIR-V IL (Platform version %4.2f)", mOclVersion); mInternals->program = clCreateProgramWithIL(mInternals->context, _binary_GPUReconstructionOCLCode_spirv_start, _binary_GPUReconstructionOCLCode_spirv_len, &ocl_error); - ocl_flags = ""; + oclBuildFlags = ""; } else #endif // clang-format on { - GPUInfo("Compiling OpenCL program from sources (Platform version %4.2f)", ver); + GPUInfo("Compiling OpenCL program from sources (Platform version %4.2f)", mOclVersion); size_t program_sizes[1] = {_binary_GPUReconstructionOCLCode_src_len}; char* programs_sources[1] = {_binary_GPUReconstructionOCLCode_src_start}; mInternals->program = clCreateProgramWithSource(mInternals->context, (cl_uint)1, (const char**)&programs_sources, program_sizes, &ocl_error); @@ -609,7 +593,7 @@ int32_t GPUReconstructionOCLBackend::GetOCLPrograms() return 1; } - if (GPUFailedMsgI(clBuildProgram(mInternals->program, 1, &mInternals->device, ocl_flags, nullptr, nullptr))) { + if (GPUFailedMsgI(clBuildProgram(mInternals->program, 1, &mInternals->device, oclBuildFlags, nullptr, nullptr))) { cl_build_status status; if (GPUFailedMsgI(clGetProgramBuildInfo(mInternals->program, mInternals->device, CL_PROGRAM_BUILD_STATUS, sizeof(status), &status, nullptr)) == 0 && status == CL_BUILD_ERROR) { size_t log_size; diff --git a/GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.h b/GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.h index d052ba53dfac5..fadb393277758 100644 --- a/GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.h +++ b/GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.h @@ -55,8 +55,6 @@ class GPUReconstructionOCLBackend : public GPUReconstructionDeviceBase void ReleaseEvent(deviceEvent ev) override; void RecordMarker(deviceEvent* ev, int32_t stream) override; - virtual bool ContextForAllPlatforms() { return false; } - template int32_t AddKernel(bool multi = false); template @@ -67,6 +65,7 @@ class GPUReconstructionOCLBackend : public GPUReconstructionDeviceBase gpu_reconstruction_kernels::krnlProperties getKernelPropertiesBackend(); GPUReconstructionOCLInternals* mInternals; + float mOclVersion; template int32_t runKernelBackend(const krnlSetupArgs& args); diff --git a/GPU/GPUTracking/Base/opencl/GPUReconstructionOCLInternals.h b/GPU/GPUTracking/Base/opencl/GPUReconstructionOCLInternals.h index 1020ce85563c2..b47c612b192d7 100644 --- a/GPU/GPUTracking/Base/opencl/GPUReconstructionOCLInternals.h +++ b/GPU/GPUTracking/Base/opencl/GPUReconstructionOCLInternals.h @@ -160,8 +160,6 @@ static inline int64_t clExecuteKernelA(cl_command_queue queue, cl_kernel krnl, s struct GPUReconstructionOCLInternals { cl_platform_id platform; cl_device_id device; - std::unique_ptr platforms; - std::unique_ptr devices; cl_context context; cl_command_queue command_queue[GPUCA_MAX_STREAMS]; cl_mem mem_gpu; diff --git a/GPU/GPUTracking/Definitions/GPUSettingsList.h b/GPU/GPUTracking/Definitions/GPUSettingsList.h index 98e0be1bdb1e5..25419f3483dd6 100644 --- a/GPU/GPUTracking/Definitions/GPUSettingsList.h +++ b/GPU/GPUTracking/Definitions/GPUSettingsList.h @@ -223,7 +223,6 @@ AddHelp("help", 'h') EndConfig() BeginSubConfig(GPUSettingsProcessing, proc, configStandalone, "PROC", 0, "Processing settings", proc) -AddOption(platformNum, int32_t, -1, "", 0, "Platform to use, in case the backend provides multiple platforms (OpenCL only, -1 = auto-select)") AddOption(deviceNum, int32_t, -1, "gpuDevice", 0, "Set GPU device to use (-1: automatic, -2: for round-robin usage in timeslice-pipeline)") AddOption(gpuDeviceOnly, bool, false, "", 0, "Use only GPU as device (i.e. no CPU for OpenCL)") AddOption(globalInitMutex, bool, false, "", 0, "Use global mutex to synchronize initialization of multiple GPU instances") @@ -296,7 +295,9 @@ AddOption(tpcApplyDebugClusterFilter, bool, false, "", 0, "Apply custom cluster AddOption(RTCcacheFolder, std::string, "./rtccache/", "", 0, "Folder in which the cache file is stored") AddOption(RTCprependCommand, std::string, "", "", 0, "Prepend RTC compilation commands by this string") AddOption(RTCoverrideArchitecture, std::string, "", "", 0, "Override arhcitecture part of RTC compilation command line") +AddOption(oclPlatformNum, int32_t, -1, "", 0, "Platform to use, in case the backend provides multiple platforms (OpenCL only, -1 = auto-select, -2 query all platforms (also incompatible))") AddOption(oclCompileFromSources, bool, false, "", 0, "Compile OpenCL binary from included source code instead of using included spirv code") +AddOption(oclOverrideSourceBuildFlags, std::string, "", "", 0, "Override OCL build flags for compilation from source, put a space for empty options") AddOption(printSettings, bool, false, "", 0, "Print all settings when initializing") AddVariable(eventDisplay, o2::gpu::GPUDisplayFrontendInterface*, nullptr) AddSubConfig(GPUSettingsProcessingRTC, rtc) From 24a8ec456ff938ea7181b1727729e0cc16b93511 Mon Sep 17 00:00:00 2001 From: David Rohr Date: Mon, 24 Feb 2025 23:05:10 +0100 Subject: [PATCH 3/3] GPU OpenCL: Make SPIR-V version settable in CMake --- GPU/GPUTracking/Base/opencl/CMakeLists.txt | 8 ++++++-- GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.cxx | 4 ++-- 2 files changed, 8 insertions(+), 4 deletions(-) diff --git a/GPU/GPUTracking/Base/opencl/CMakeLists.txt b/GPU/GPUTracking/Base/opencl/CMakeLists.txt index 2a361356283a8..d6aa945fc77b7 100644 --- a/GPU/GPUTracking/Base/opencl/CMakeLists.txt +++ b/GPU/GPUTracking/Base/opencl/CMakeLists.txt @@ -39,6 +39,10 @@ set(OCL_DEFINECL "-D$= 17 @@ -47,7 +51,7 @@ if(OPENCL_ENABLED_SPIRV) # BUILD OpenCL intermediate code for SPIR-V target OUTPUT ${CL_BIN}.spirv COMMAND ${CMAKE_COMMAND} -E env "PATH=${TMP_LLVM_SPIRV_PATH}:\$$PATH" ${LLVM_CLANG} -O0 - --target=spirv64 + --target=spirv64v${GPUCA_OCL_SPIRV_VERSION} -fno-integrated-objemitter -ferror-limit=1000 -Wno-invalid-constexpr -Wno-unused-command-line-argument ${OCL_FLAGS} @@ -106,4 +110,4 @@ endif() if(OPENCL_ENABLED_SPIRV) target_compile_definitions(${targetName} PRIVATE OPENCL_ENABLED_SPIRV) endif() -target_compile_definitions(${targetName} PRIVATE GPUCA_OCL_BUILD_FLAGS=$) +target_compile_definitions(${targetName} PRIVATE GPUCA_OCL_BUILD_FLAGS=$ GPUCA_OCL_SPIRV_VERSION=${GPUCA_OCL_SPIRV_VERSION}) diff --git a/GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.cxx b/GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.cxx index 6b918fe501330..5fac46a214815 100644 --- a/GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.cxx +++ b/GPU/GPUTracking/Base/opencl/GPUReconstructionOCL.cxx @@ -209,14 +209,14 @@ int32_t GPUReconstructionOCLBackend::InitDevice_Runtime() while ((pos = device_il_version.find("SPIR-V", pos)) != std::string::npos) { float spirvVersion; sscanf(device_il_version.c_str() + pos, "SPIR-V_%f", &spirvVersion); - if (spirvVersion >= 1.2) { + if (spirvVersion >= GPUCA_OCL_SPIRV_VERSION) { break; } pos += strlen("SPIR-V_0.0"); } if (pos == std::string::npos) { deviceOK = false; - deviceFailure += " - No SPIR-V 1.6 (" + device_il_version + ")"; + deviceFailure += " - No SPIR-V " + std::to_string(GPUCA_OCL_SPIRV_VERSION) + " (" + device_il_version + ")"; } }