From 430cca3f4ce830494945833d7b2024801fe465ed Mon Sep 17 00:00:00 2001 From: Mateusz Jablonski Date: Wed, 28 Oct 2020 17:18:37 +0100 Subject: [PATCH] Use all passed devices when linking program Related-To: NEO-5001 Change-Id: I3944375023b9cb3322df83e6e042ba8be2bd762c Signed-off-by: Mateusz Jablonski --- opencl/source/api/api.cpp | 47 ++--- opencl/source/cl_device/cl_device.cpp | 2 +- opencl/source/cl_device/cl_device_vector.h | 2 +- opencl/source/program/link.cpp | 100 +++++----- opencl/source/program/program.cpp | 24 +++ opencl/source/program/program.h | 4 +- .../api/cl_get_program_build_info_tests.inl | 175 +++++++++++++++++- .../unit_test/api/cl_link_program_tests.inl | 108 ++++++++++- .../test/unit_test/program/program_tests.cpp | 48 ++--- .../program_with_block_kernels_tests.cpp | 2 +- .../program_with_kernel_debug_tests.cpp | 4 +- 11 files changed, 397 insertions(+), 119 deletions(-) diff --git a/opencl/source/api/api.cpp b/opencl/source/api/api.cpp index cbadda4e2b..fc54a579bc 100644 --- a/opencl/source/api/api.cpp +++ b/opencl/source/api/api.cpp @@ -1539,30 +1539,10 @@ cl_int CL_API_CALL clCompileProgram(cl_program program, retVal = validateObjects(WithCastToInternal(program, &pProgram), Program::isValidCallback(funcNotify, userData)); ClDeviceVector deviceVector; - const ClDeviceVector *deviceVectorPtr = &deviceVector; + ClDeviceVector *deviceVectorPtr = &deviceVector; if (CL_SUCCESS == retVal) { - if (deviceList == nullptr) { - if (numDevices == 0) { - deviceVectorPtr = &pProgram->getDevices(); - } else { - retVal = CL_INVALID_VALUE; - } - - } else { - if (numDevices == 0) { - retVal = CL_INVALID_VALUE; - } else { - for (auto i = 0u; i < numDevices; i++) { - auto device = castToObject(deviceList[i]); - if (!device || !pProgram->isDeviceAssociated(*device)) { - retVal = CL_INVALID_DEVICE; - break; - } - deviceVector.push_back(device); - } - } - } + retVal = Program::processInputDevices(deviceVectorPtr, numDevices, deviceList, pProgram->getDevices()); } if (CL_SUCCESS == retVal) { retVal = pProgram->compile(*deviceVectorPtr, options, @@ -1590,23 +1570,28 @@ cl_program CL_API_CALL clLinkProgram(cl_context context, ErrorCodeHelper err(errcodeRet, CL_SUCCESS); Context *pContext = nullptr; - Program *program = nullptr; + Program *pProgram = nullptr; retVal = validateObjects(WithCastToInternal(context, &pContext), Program::isValidCallback(funcNotify, userData)); + + ClDeviceVector deviceVector; + ClDeviceVector *deviceVectorPtr = &deviceVector; + if (CL_SUCCESS == retVal) { + retVal = Program::processInputDevices(deviceVectorPtr, numDevices, deviceList, pContext->getDevices()); + } + if (CL_SUCCESS == retVal) { - ClDeviceVector deviceVector; - deviceVector.push_back(pContext->getDevice(0)); - program = new Program(pContext, false, deviceVector); - retVal = program->link(numDevices, deviceList, options, - numInputPrograms, inputPrograms); - program->invokeCallback(funcNotify, userData); + pProgram = new Program(pContext, false, *deviceVectorPtr); + retVal = pProgram->link(*deviceVectorPtr, options, + numInputPrograms, inputPrograms); + pProgram->invokeCallback(funcNotify, userData); } err.set(retVal); - TRACING_EXIT(clLinkProgram, (cl_program *)&program); - return program; + TRACING_EXIT(clLinkProgram, (cl_program *)&pProgram); + return pProgram; } cl_int CL_API_CALL clUnloadPlatformCompiler(cl_platform_id platform) { diff --git a/opencl/source/cl_device/cl_device.cpp b/opencl/source/cl_device/cl_device.cpp index 0009231daf..c73d28c0c1 100644 --- a/opencl/source/cl_device/cl_device.cpp +++ b/opencl/source/cl_device/cl_device.cpp @@ -169,7 +169,7 @@ ClDeviceVector::ClDeviceVector(const cl_device_id *devices, } } -void ClDeviceVector::toDeviceIDs(std::vector &devIDs) { +void ClDeviceVector::toDeviceIDs(std::vector &devIDs) const { int i = 0; devIDs.resize(this->size()); diff --git a/opencl/source/cl_device/cl_device_vector.h b/opencl/source/cl_device/cl_device_vector.h index 70867559cc..c3f1a5a0f3 100644 --- a/opencl/source/cl_device/cl_device_vector.h +++ b/opencl/source/cl_device/cl_device_vector.h @@ -20,7 +20,7 @@ class ClDeviceVector : public StackVec { ClDeviceVector &operator=(const ClDeviceVector &) = default; ClDeviceVector(const cl_device_id *devices, cl_uint numDevices); - void toDeviceIDs(std::vector &devIDs); + void toDeviceIDs(std::vector &devIDs) const; }; } // namespace NEO diff --git a/opencl/source/program/link.cpp b/opencl/source/program/link.cpp index 5421523001..afa97c8708 100644 --- a/opencl/source/program/link.cpp +++ b/opencl/source/program/link.cpp @@ -28,36 +28,30 @@ namespace NEO { cl_int Program::link( - cl_uint numDevices, - const cl_device_id *deviceList, + const ClDeviceVector &deviceVector, const char *buildOptions, cl_uint numInputPrograms, const cl_program *inputPrograms) { cl_int retVal = CL_SUCCESS; bool isCreateLibrary; - auto clDevice = this->pDevice->getSpecializedDevice(); - UNRECOVERABLE_IF(clDevice == nullptr); + auto defaultClDevice = deviceVector[0]; + UNRECOVERABLE_IF(defaultClDevice == nullptr); + auto &defaultDevice = defaultClDevice->getDevice(); + internalOptions.clear(); do { - if (((deviceList == nullptr) && (numDevices != 0)) || - ((deviceList != nullptr) && (numDevices == 0))) { - retVal = CL_INVALID_VALUE; - break; - } - if ((numInputPrograms == 0) || (inputPrograms == nullptr)) { retVal = CL_INVALID_VALUE; break; } - if ((deviceList != nullptr) && validateObject(*deviceList) != CL_SUCCESS) { - retVal = CL_INVALID_DEVICE; + if (std::any_of(deviceVector.begin(), deviceVector.end(), [&](auto device) { return CL_BUILD_IN_PROGRESS == buildStatuses[device]; })) { + retVal = CL_INVALID_OPERATION; break; } - if (buildStatuses[clDevice] == CL_BUILD_IN_PROGRESS) { - retVal = CL_INVALID_OPERATION; - break; + for (const auto &device : deviceVector) { + buildStatuses[device] = CL_BUILD_IN_PROGRESS; } options = (buildOptions != nullptr) ? buildOptions : ""; @@ -76,8 +70,6 @@ cl_int Program::link( isCreateLibrary = CompilerOptions::contains(options, CompilerOptions::createLibrary); - buildStatuses[clDevice] = CL_BUILD_IN_PROGRESS; - NEO::Elf::ElfEncoder<> elfEncoder(true, false, 1U); elfEncoder.getElfFileHeader().type = NEO::Elf::ET_OPENCL_OBJECTS; @@ -126,7 +118,7 @@ cl_int Program::link( auto clLinkInput = elfEncoder.encode(); - CompilerInterface *pCompilerInterface = pDevice->getCompilerInterface(); + CompilerInterface *pCompilerInterface = defaultDevice.getCompilerInterface(); if (!pCompilerInterface) { retVal = CL_OUT_OF_HOST_MEMORY; break; @@ -140,43 +132,46 @@ cl_int Program::link( inputArgs.GTPinInput = gtpinGetIgcInit(); if (!isCreateLibrary) { - inputArgs.outType = IGC::CodeType::oclGenBin; - NEO::TranslationOutput compilerOuput = {}; - auto compilerErr = pCompilerInterface->link(this->getDevice(), inputArgs, compilerOuput); - this->updateBuildLog(this->pDevice->getRootDeviceIndex(), compilerOuput.frontendCompilerLog.c_str(), compilerOuput.frontendCompilerLog.size()); - this->updateBuildLog(this->pDevice->getRootDeviceIndex(), compilerOuput.backendCompilerLog.c_str(), compilerOuput.backendCompilerLog.size()); - retVal = asClError(compilerErr); - if (retVal != CL_SUCCESS) { - break; - } + for (const auto &device : deviceVector) { + inputArgs.outType = IGC::CodeType::oclGenBin; + NEO::TranslationOutput compilerOuput = {}; + auto compilerErr = pCompilerInterface->link(device->getDevice(), inputArgs, compilerOuput); + this->updateBuildLog(device->getRootDeviceIndex(), compilerOuput.frontendCompilerLog.c_str(), compilerOuput.frontendCompilerLog.size()); + this->updateBuildLog(device->getRootDeviceIndex(), compilerOuput.backendCompilerLog.c_str(), compilerOuput.backendCompilerLog.size()); + retVal = asClError(compilerErr); + if (retVal != CL_SUCCESS) { + break; + } - this->replaceDeviceBinary(std::move(compilerOuput.deviceBinary.mem), compilerOuput.deviceBinary.size, pDevice->getRootDeviceIndex()); - this->debugData = std::move(compilerOuput.debugData.mem); - this->debugDataSize = compilerOuput.debugData.size; + this->replaceDeviceBinary(std::move(compilerOuput.deviceBinary.mem), compilerOuput.deviceBinary.size, device->getRootDeviceIndex()); + this->debugData = std::move(compilerOuput.debugData.mem); + this->debugDataSize = compilerOuput.debugData.size; - retVal = processGenBinary(pDevice->getRootDeviceIndex()); - if (retVal != CL_SUCCESS) { - break; - } - programBinaryType = CL_PROGRAM_BINARY_TYPE_EXECUTABLE; + retVal = processGenBinary(device->getRootDeviceIndex()); + if (retVal != CL_SUCCESS) { + break; + } + programBinaryType = CL_PROGRAM_BINARY_TYPE_EXECUTABLE; - if (isKernelDebugEnabled()) { - processDebugData(); - auto clDevice = this->getDevice().getSpecializedDevice(); - UNRECOVERABLE_IF(clDevice == nullptr); - for (auto kernelInfo : kernelInfoArray) { - clDevice->getSourceLevelDebugger()->notifyKernelDebugData(&kernelInfo->debugData, - kernelInfo->kernelDescriptor.kernelMetadata.kernelName, - kernelInfo->heapInfo.pKernelHeap, - kernelInfo->heapInfo.KernelHeapSize); + if (isKernelDebugEnabled()) { + processDebugData(); + for (auto kernelInfo : kernelInfoArray) { + device->getSourceLevelDebugger()->notifyKernelDebugData(&kernelInfo->debugData, + kernelInfo->kernelDescriptor.kernelMetadata.kernelName, + kernelInfo->heapInfo.pKernelHeap, + kernelInfo->heapInfo.KernelHeapSize); + } } } + } else { inputArgs.outType = IGC::CodeType::llvmBc; NEO::TranslationOutput compilerOuput = {}; - auto compilerErr = pCompilerInterface->createLibrary(*this->pDevice, inputArgs, compilerOuput); - this->updateBuildLog(this->pDevice->getRootDeviceIndex(), compilerOuput.frontendCompilerLog.c_str(), compilerOuput.frontendCompilerLog.size()); - this->updateBuildLog(this->pDevice->getRootDeviceIndex(), compilerOuput.backendCompilerLog.c_str(), compilerOuput.backendCompilerLog.size()); + auto compilerErr = pCompilerInterface->createLibrary(defaultDevice, inputArgs, compilerOuput); + for (const auto &device : deviceVector) { + this->updateBuildLog(device->getRootDeviceIndex(), compilerOuput.frontendCompilerLog.c_str(), compilerOuput.frontendCompilerLog.size()); + this->updateBuildLog(device->getRootDeviceIndex(), compilerOuput.backendCompilerLog.c_str(), compilerOuput.backendCompilerLog.size()); + } retVal = asClError(compilerErr); if (retVal != CL_SUCCESS) { break; @@ -188,15 +183,22 @@ cl_int Program::link( this->debugDataSize = compilerOuput.debugData.size; programBinaryType = CL_PROGRAM_BINARY_TYPE_LIBRARY; } + if (retVal != CL_SUCCESS) { + break; + } updateNonUniformFlag(&*inputProgramsInternal.begin(), inputProgramsInternal.size()); separateBlockKernels(); } while (false); if (retVal != CL_SUCCESS) { - buildStatuses[clDevice] = CL_BUILD_ERROR; + for (const auto &device : deviceVector) { + buildStatuses[device] = CL_BUILD_ERROR; + } programBinaryType = CL_PROGRAM_BINARY_TYPE_NONE; } else { - buildStatuses[clDevice] = CL_BUILD_SUCCESS; + for (const auto &device : deviceVector) { + buildStatuses[device] = CL_BUILD_SUCCESS; + } } internalOptions.clear(); diff --git a/opencl/source/program/program.cpp b/opencl/source/program/program.cpp index 6cc89cd94c..db97948fa7 100644 --- a/opencl/source/program/program.cpp +++ b/opencl/source/program/program.cpp @@ -502,4 +502,28 @@ void Program::invokeCallback(void(CL_CALLBACK *funcNotify)(cl_program program, v bool Program::isDeviceAssociated(const ClDevice &clDevice) const { return std::any_of(clDevices.begin(), clDevices.end(), [&](auto programDevice) { return programDevice == &clDevice; }); } + +cl_int Program::processInputDevices(ClDeviceVector *&deviceVectorPtr, cl_uint numDevices, const cl_device_id *deviceList, const ClDeviceVector &allAvailableDevices) { + if (deviceList == nullptr) { + if (numDevices == 0) { + deviceVectorPtr = const_cast(&allAvailableDevices); + } else { + return CL_INVALID_VALUE; + } + + } else { + if (numDevices == 0) { + return CL_INVALID_VALUE; + } else { + for (auto i = 0u; i < numDevices; i++) { + auto device = castToObject(deviceList[i]); + if (!device || !std::any_of(allAvailableDevices.begin(), allAvailableDevices.end(), [&](auto validDevice) { return validDevice == device; })) { + return CL_INVALID_DEVICE; + } + deviceVectorPtr->push_back(device); + } + } + } + return CL_SUCCESS; +} } // namespace NEO diff --git a/opencl/source/program/program.h b/opencl/source/program/program.h index fe9c0f1b3f..a0d1061a4c 100644 --- a/opencl/source/program/program.h +++ b/opencl/source/program/program.h @@ -138,7 +138,7 @@ class Program : public BaseObject<_cl_program> { cl_int compile(const ClDeviceVector &deviceVector, const char *buildOptions, cl_uint numInputHeaders, const cl_program *inputHeaders, const char **headerIncludeNames); - cl_int link(cl_uint numDevices, const cl_device_id *deviceList, const char *buildOptions, + cl_int link(const ClDeviceVector &deviceVector, const char *buildOptions, cl_uint numInputPrograms, const cl_program *inputPrograms); cl_int setProgramSpecializationConstant(cl_uint specId, size_t specSize, const void *specValue); @@ -266,6 +266,8 @@ class Program : public BaseObject<_cl_program> { const ClDeviceVector &getDevices() const { return clDevices; } bool isDeviceAssociated(const ClDevice &clDevice) const; + static cl_int processInputDevices(ClDeviceVector *&deviceVectorPtr, cl_uint numDevices, const cl_device_id *deviceList, const ClDeviceVector &allAvailableDevices); + protected: MOCKABLE_VIRTUAL cl_int createProgramFromBinary(const void *pBinary, size_t binarySize, uint32_t rootDeviceIndex); diff --git a/opencl/test/unit_test/api/cl_get_program_build_info_tests.inl b/opencl/test/unit_test/api/cl_get_program_build_info_tests.inl index 9885c9ff5a..a95237862e 100644 --- a/opencl/test/unit_test/api/cl_get_program_build_info_tests.inl +++ b/opencl/test/unit_test/api/cl_get_program_build_info_tests.inl @@ -20,10 +20,10 @@ #include "cl_api_tests.h" using namespace NEO; - typedef api_tests clGetProgramBuildInfoTests; namespace ULT { +void verifyDevices(cl_program pProgram, size_t expectedNumDevices, cl_device_id *expectedDevices); TEST_F(clGetProgramBuildInfoTests, givenSourceWhenclGetProgramBuildInfoIsCalledThenReturnClBuildNone) { cl_program pProgram = nullptr; @@ -278,4 +278,177 @@ TEST_F(clGetProgramBuildInfoTests, givenInvalidDeviceInputWhenGetProgramBuildInf retVal = clGetProgramBuildInfo(pProgram, context.getDevice(0), CL_PROGRAM_BUILD_STATUS, sizeof(buildStatus), &buildStatus, nullptr); EXPECT_EQ(CL_INVALID_DEVICE, retVal); } + +TEST(clGetProgramBuildInfoTest, givenMultiDeviceProgramWhenLinkingForSpecificDevicesThenOnlySpecificDevicesReportBuildStatus) { + MockProgram *pProgram = nullptr; + std::unique_ptr pSource = nullptr; + size_t sourceSize = 0; + std::string testFile; + + KernelBinaryHelper kbHelper("CopyBuffer_simd16"); + + testFile.append(clFiles); + testFile.append("CopyBuffer_simd16.cl"); + + pSource = loadDataFromFile( + testFile.c_str(), + sourceSize); + + ASSERT_NE(0u, sourceSize); + ASSERT_NE(nullptr, pSource); + + const char *sources[1] = {pSource.get()}; + + MockUnrestrictiveContextMultiGPU context; + cl_int retVal = CL_INVALID_PROGRAM; + + pProgram = Program::create( + &context, + 1, + sources, + &sourceSize, + retVal); + + EXPECT_NE(nullptr, pProgram); + ASSERT_EQ(CL_SUCCESS, retVal); + + cl_build_status buildStatus; + for (const auto &device : context.getDevices()) { + retVal = clGetProgramBuildInfo(pProgram, device, CL_PROGRAM_BUILD_STATUS, sizeof(buildStatus), &buildStatus, NULL); + EXPECT_EQ(CL_SUCCESS, retVal); + EXPECT_EQ(CL_BUILD_NONE, buildStatus); + } + + retVal = clCompileProgram( + pProgram, + 0, + nullptr, + nullptr, + 0, + nullptr, + nullptr, + nullptr, + nullptr); + + ASSERT_EQ(CL_SUCCESS, retVal); + + pProgram->setBuildStatus(CL_BUILD_NONE); + + cl_device_id devicesForLinking[] = {context.getDevice(1), context.getDevice(3)}; + cl_program programForLinking = pProgram; + + auto outProgram = clLinkProgram( + &context, + 2, + devicesForLinking, + nullptr, + 1, + &programForLinking, + nullptr, + nullptr, + &retVal); + + ASSERT_EQ(CL_SUCCESS, retVal); + EXPECT_NE(nullptr, outProgram); + + verifyDevices(outProgram, 2, devicesForLinking); + + for (const auto &device : devicesForLinking) { + retVal = clGetProgramBuildInfo(outProgram, device, CL_PROGRAM_BUILD_STATUS, sizeof(buildStatus), &buildStatus, NULL); + EXPECT_EQ(CL_SUCCESS, retVal); + EXPECT_EQ(CL_BUILD_SUCCESS, buildStatus); + } + + retVal = clReleaseProgram(pProgram); + EXPECT_EQ(CL_SUCCESS, retVal); + retVal = clReleaseProgram(outProgram); + EXPECT_EQ(CL_SUCCESS, retVal); +} + +TEST(clGetProgramBuildInfoTest, givenMultiDeviceProgramWhenLinkingWithoutInputDevicesThenAllDevicesReportBuildStatus) { + MockProgram *pProgram = nullptr; + std::unique_ptr pSource = nullptr; + size_t sourceSize = 0; + std::string testFile; + + KernelBinaryHelper kbHelper("CopyBuffer_simd16"); + + testFile.append(clFiles); + testFile.append("CopyBuffer_simd16.cl"); + + pSource = loadDataFromFile( + testFile.c_str(), + sourceSize); + + ASSERT_NE(0u, sourceSize); + ASSERT_NE(nullptr, pSource); + + const char *sources[1] = {pSource.get()}; + + MockUnrestrictiveContextMultiGPU context; + cl_int retVal = CL_INVALID_PROGRAM; + + pProgram = Program::create( + &context, + 1, + sources, + &sourceSize, + retVal); + + EXPECT_NE(nullptr, pProgram); + ASSERT_EQ(CL_SUCCESS, retVal); + + cl_build_status buildStatus; + for (const auto &device : context.getDevices()) { + retVal = clGetProgramBuildInfo(pProgram, device, CL_PROGRAM_BUILD_STATUS, sizeof(buildStatus), &buildStatus, NULL); + EXPECT_EQ(CL_SUCCESS, retVal); + EXPECT_EQ(CL_BUILD_NONE, buildStatus); + } + + retVal = clCompileProgram( + pProgram, + 0, + nullptr, + nullptr, + 0, + nullptr, + nullptr, + nullptr, + nullptr); + + ASSERT_EQ(CL_SUCCESS, retVal); + + pProgram->setBuildStatus(CL_BUILD_NONE); + + cl_program programForLinking = pProgram; + + auto outProgram = clLinkProgram( + &context, + 0, + nullptr, + nullptr, + 1, + &programForLinking, + nullptr, + nullptr, + &retVal); + + ASSERT_EQ(CL_SUCCESS, retVal); + EXPECT_NE(nullptr, outProgram); + + std::vector contextDevices; + context.getDevices().toDeviceIDs(contextDevices); + verifyDevices(outProgram, contextDevices.size(), contextDevices.data()); + + for (const auto &device : context.getDevices()) { + retVal = clGetProgramBuildInfo(outProgram, device, CL_PROGRAM_BUILD_STATUS, sizeof(buildStatus), &buildStatus, NULL); + EXPECT_EQ(CL_SUCCESS, retVal); + EXPECT_EQ(CL_BUILD_SUCCESS, buildStatus); + } + + retVal = clReleaseProgram(pProgram); + EXPECT_EQ(CL_SUCCESS, retVal); + retVal = clReleaseProgram(outProgram); + EXPECT_EQ(CL_SUCCESS, retVal); +} } // namespace ULT diff --git a/opencl/test/unit_test/api/cl_link_program_tests.inl b/opencl/test/unit_test/api/cl_link_program_tests.inl index 31336465f0..1cc52f25c5 100644 --- a/opencl/test/unit_test/api/cl_link_program_tests.inl +++ b/opencl/test/unit_test/api/cl_link_program_tests.inl @@ -206,7 +206,7 @@ TEST_F(clLinkProgramTests, GivenProgramsWithSpecConstantsThenSpecConstantsAreEmb MockCompilerDebugVars igcDebugVars; igcDebugVars.receivedInput = &receivedInput; gEnvironment->igcPushDebugVars(igcDebugVars); - progDst->link(0U, nullptr, "", 3, inputPrograms); + progDst->link(progDst->getDevices(), "", 3, inputPrograms); gEnvironment->igcPopDebugVars(); std::string elfDecodeError; @@ -374,4 +374,110 @@ TEST_F(clLinkProgramTests, GivenValidCallbackInputWhenLinkProgramThenCallbackIsI EXPECT_EQ(CL_SUCCESS, retVal); } +TEST_F(clLinkProgramTests, givenMultiDeviceProgramWhenLinkingForInvalidDevicesInputThenInvalidDeviceErrorIsReturned) { + cl_program pProgram = nullptr; + size_t sourceSize = 0; + std::string testFile; + + testFile.append(clFiles); + testFile.append("copybuffer.cl"); + auto pSource = loadDataFromFile( + testFile.c_str(), + sourceSize); + + ASSERT_NE(0u, sourceSize); + ASSERT_NE(nullptr, pSource); + + const char *sources[1] = {pSource.get()}; + pProgram = clCreateProgramWithSource( + pContext, + 1, + sources, + &sourceSize, + &retVal); + + EXPECT_NE(nullptr, pProgram); + ASSERT_EQ(CL_SUCCESS, retVal); + + retVal = clCompileProgram( + pProgram, + 1, + &testedClDevice, + nullptr, + 0, + nullptr, + nullptr, + nullptr, + nullptr); + + ASSERT_EQ(CL_SUCCESS, retVal); + + cl_program program = pProgram; + cl_program outProgram; + + MockContext mockContext; + cl_device_id nullDeviceInput[] = {pContext->getDevice(0), nullptr}; + cl_device_id notAssociatedDeviceInput[] = {mockContext.getDevice(0)}; + cl_device_id validDeviceInput[] = {pContext->getDevice(0)}; + + outProgram = clLinkProgram( + pContext, + 0, + validDeviceInput, + nullptr, + 1, + &program, + nullptr, + nullptr, + &retVal); + + EXPECT_EQ(CL_INVALID_VALUE, retVal); + EXPECT_EQ(nullptr, outProgram); + + outProgram = clLinkProgram( + pContext, + 1, + nullptr, + nullptr, + 1, + &program, + nullptr, + nullptr, + &retVal); + + EXPECT_EQ(CL_INVALID_VALUE, retVal); + EXPECT_EQ(nullptr, outProgram); + + outProgram = clLinkProgram( + pContext, + 2, + nullDeviceInput, + nullptr, + 1, + &program, + nullptr, + nullptr, + &retVal); + + EXPECT_EQ(CL_INVALID_DEVICE, retVal); + EXPECT_EQ(nullptr, outProgram); + + outProgram = clLinkProgram( + pContext, + 1, + notAssociatedDeviceInput, + nullptr, + 1, + &program, + nullptr, + nullptr, + &retVal); + + EXPECT_EQ(CL_INVALID_DEVICE, retVal); + EXPECT_EQ(nullptr, outProgram); + + retVal = clReleaseProgram(pProgram); + EXPECT_EQ(CL_SUCCESS, retVal); +} + } // namespace ULT diff --git a/opencl/test/unit_test/program/program_tests.cpp b/opencl/test/unit_test/program/program_tests.cpp index 639a86f678..a85737f038 100644 --- a/opencl/test/unit_test/program/program_tests.cpp +++ b/opencl/test/unit_test/program/program_tests.cpp @@ -1113,7 +1113,7 @@ TEST_F(ProgramTests, GivenFlagsWhenLinkingProgramThenBuildOptionsHaveBeenApplied pDevice->getExecutionEnvironment()->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->compilerInterface.reset(cip); - retVal = pProgram->link(0, nullptr, options.c_str(), 1, &program); + retVal = pProgram->link(pProgram->getDevices(), options.c_str(), 1, &program); EXPECT_EQ(CL_SUCCESS, retVal); // Check build options that were applied @@ -1182,7 +1182,6 @@ TEST_P(ProgramFromSourceTest, GivenSpecificParamatersWhenLinkingProgramThenSucce &usedDevice, SourceFileName); - cl_device_id deviceList = {0}; cl_program program = pProgram; cl_program nullprogram = nullptr; cl_program invprogram = (cl_program)pContext; @@ -1190,36 +1189,25 @@ TEST_P(ProgramFromSourceTest, GivenSpecificParamatersWhenLinkingProgramThenSucce // Order of following microtests is important - do not change. // Add new microtests at end. - // invalid link parameters: combinations of numDevices & deviceList - retVal = pProgram->link(1, nullptr, nullptr, 1, &program); - EXPECT_EQ(CL_INVALID_VALUE, retVal); - - retVal = pProgram->link(0, &deviceList, nullptr, 1, &program); - EXPECT_EQ(CL_INVALID_VALUE, retVal); - // invalid link parameters: combinations of numInputPrograms & inputPrograms - retVal = pProgram->link(0, nullptr, nullptr, 0, &program); + retVal = pProgram->link(pProgram->getDevices(), nullptr, 0, &program); EXPECT_EQ(CL_INVALID_VALUE, retVal); - retVal = pProgram->link(0, nullptr, nullptr, 1, nullptr); + retVal = pProgram->link(pProgram->getDevices(), nullptr, 1, nullptr); EXPECT_EQ(CL_INVALID_VALUE, retVal); - // invalid link parameters: invalid content of deviceList - retVal = pProgram->link(1, &deviceList, nullptr, 1, &program); - EXPECT_EQ(CL_INVALID_DEVICE, retVal); - // fail linking - another linking is already in progress pProgram->setBuildStatus(CL_BUILD_IN_PROGRESS); - retVal = pProgram->link(0, nullptr, nullptr, 1, &program); + retVal = pProgram->link(pProgram->getDevices(), nullptr, 1, &program); EXPECT_EQ(CL_INVALID_OPERATION, retVal); pProgram->setBuildStatus(CL_BUILD_NONE); // invalid link parameters: invalid Program object==nullptr - retVal = pProgram->link(0, nullptr, nullptr, 1, &nullprogram); + retVal = pProgram->link(pProgram->getDevices(), nullptr, 1, &nullprogram); EXPECT_EQ(CL_INVALID_PROGRAM, retVal); // invalid link parameters: invalid Program object==non Program object - retVal = pProgram->link(0, nullptr, nullptr, 1, &invprogram); + retVal = pProgram->link(pProgram->getDevices(), nullptr, 1, &invprogram); EXPECT_EQ(CL_INVALID_PROGRAM, retVal); // compile successfully a kernel to be linked later @@ -1232,29 +1220,29 @@ TEST_P(ProgramFromSourceTest, GivenSpecificParamatersWhenLinkingProgramThenSucce pProgram->irBinary.release(); size_t irBinSize = pProgram->irBinarySize; pProgram->setIrBinary(nullptr, false); - retVal = pProgram->link(0, nullptr, nullptr, 1, &program); + retVal = pProgram->link(pProgram->getDevices(), nullptr, 1, &program); EXPECT_EQ(CL_INVALID_PROGRAM, retVal); pProgram->setIrBinary(pIrBin, isSpirvTmp); // fail linking - size of code to be linked is == 0 pProgram->setIrBinarySize(0, isSpirvTmp); - retVal = pProgram->link(0, nullptr, nullptr, 1, &program); + retVal = pProgram->link(pProgram->getDevices(), nullptr, 1, &program); EXPECT_EQ(CL_INVALID_PROGRAM, retVal); pProgram->setIrBinarySize(irBinSize, isSpirvTmp); // fail linking - any link error (here caused by specifying unrecognized option) - retVal = pProgram->link(0, nullptr, "-invalid-option", 1, &program); + retVal = pProgram->link(pProgram->getDevices(), "-invalid-option", 1, &program); EXPECT_EQ(CL_LINK_PROGRAM_FAILURE, retVal); // fail linking - linked code is corrupted and cannot be postprocessed auto device = static_cast(usedDevice); auto p2 = std::make_unique(toClDeviceVector(*device)); - retVal = p2->link(0, nullptr, nullptr, 1, &program); + retVal = p2->link(p2->getDevices(), nullptr, 1, &program); EXPECT_EQ(CL_INVALID_BINARY, retVal); p2.reset(nullptr); // link successfully - retVal = pProgram->link(0, nullptr, nullptr, 1, &program); + retVal = pProgram->link(pProgram->getDevices(), nullptr, 1, &program); EXPECT_EQ(CL_SUCCESS, retVal); } @@ -1269,11 +1257,11 @@ TEST_P(ProgramFromSourceTest, GivenInvalidOptionsWhenCreatingLibraryThenCorrectE EXPECT_EQ(CL_SUCCESS, retVal); // create library successfully - retVal = pProgram->link(0, nullptr, CompilerOptions::createLibrary.data(), 1, &program); + retVal = pProgram->link(pProgram->getDevices(), CompilerOptions::createLibrary.data(), 1, &program); EXPECT_EQ(CL_SUCCESS, retVal); // fail library creation - any link error (here caused by specifying unrecognized option) - retVal = pProgram->link(0, nullptr, CompilerOptions::concatenate(CompilerOptions::createLibrary, "-invalid-option").c_str(), 1, &program); + retVal = pProgram->link(pProgram->getDevices(), CompilerOptions::concatenate(CompilerOptions::createLibrary, "-invalid-option").c_str(), 1, &program); EXPECT_EQ(CL_LINK_PROGRAM_FAILURE, retVal); auto device = pContext->getDevice(0); @@ -1283,7 +1271,7 @@ TEST_P(ProgramFromSourceTest, GivenInvalidOptionsWhenCreatingLibraryThenCorrectE auto failingProgram = std::make_unique(toClDeviceVector(*device)); // fail library creation - CompilerInterface cannot be obtained - retVal = failingProgram->link(0, nullptr, CompilerOptions::createLibrary.data(), 1, &program); + retVal = failingProgram->link(failingProgram->getDevices(), CompilerOptions::createLibrary.data(), 1, &program); EXPECT_EQ(CL_OUT_OF_HOST_MEMORY, retVal); std::swap(rootDeviceEnvironment, executionEnvironment->rootDeviceEnvironments[device->getRootDeviceIndex()]); } @@ -2335,7 +2323,7 @@ TEST_F(ProgramTests, WhenLinkingTwoValidSpirvProgramsThenValidProgramIsReturned) EXPECT_EQ(CL_SUCCESS, errCode); cl_program linkNodes[] = {node1, node2}; - errCode = prog->link(0, nullptr, nullptr, 2, linkNodes); + errCode = prog->link(prog->getDevices(), nullptr, 2, linkNodes); EXPECT_EQ(CL_SUCCESS, errCode); prog->release(); @@ -2886,7 +2874,6 @@ TEST_F(ProgramBinTest, GivenBuildWithDebugDataThenBuildDataAvailableViaGetInfo) TEST_F(ProgramBinTest, GivenDebugDataAvailableWhenLinkingProgramThenDebugDataIsStoredInProgram) { DebugDataGuard debugDataGuard; - cl_device_id device = pClDevice; const char *sourceCode = "__kernel void\nCB(\n__global unsigned int* src, __global unsigned int* dst)\n{\nint id = (int)get_global_id(0);\ndst[id] = src[id];\n}\n"; pProgram = Program::create( pContext, @@ -2899,7 +2886,7 @@ TEST_F(ProgramBinTest, GivenDebugDataAvailableWhenLinkingProgramThenDebugDataIsS EXPECT_EQ(CL_SUCCESS, retVal); cl_program programToLink = pProgram; - retVal = pProgram->link(1, &device, nullptr, 1, &programToLink); + retVal = pProgram->link(pProgram->getDevices(), nullptr, 1, &programToLink); EXPECT_EQ(CL_SUCCESS, retVal); EXPECT_NE(nullptr, pProgram->getDebugData()); @@ -2965,7 +2952,6 @@ class MockCompilerInterfaceWithGtpinParam : public CompilerInterface { }; TEST_F(ProgramBinTest, GivenSourceKernelWhenLinkingProgramThenGtpinInitInfoIsPassed) { - cl_device_id device = pClDevice; void *pIgcInitPtr = reinterpret_cast(0x1234); gtpinSetIgcInit(pIgcInitPtr); const char *sourceCode = "__kernel void\nCB(\n__global unsigned int* src, __global unsigned int* dst)\n{\nint id = (int)get_global_id(0);\ndst[id] = src[id];\n}\n"; @@ -2982,7 +2968,7 @@ TEST_F(ProgramBinTest, GivenSourceKernelWhenLinkingProgramThenGtpinInitInfoIsPas pDevice->getExecutionEnvironment()->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->compilerInterface.reset(mockCompilerInterface.get()); cl_program programToLink = pProgram; - retVal = pProgram->link(1, &device, nullptr, 1, &programToLink); + retVal = pProgram->link(pProgram->getDevices(), nullptr, 1, &programToLink); EXPECT_EQ(pIgcInitPtr, mockCompilerInterface->gtpinInfoPassed); mockCompilerInterface.release(); diff --git a/opencl/test/unit_test/program/program_with_block_kernels_tests.cpp b/opencl/test/unit_test/program/program_with_block_kernels_tests.cpp index bac4e7ea18..6dbbdffc54 100644 --- a/opencl/test/unit_test/program/program_with_block_kernels_tests.cpp +++ b/opencl/test/unit_test/program/program_with_block_kernels_tests.cpp @@ -104,7 +104,7 @@ TEST_F(ProgramWithBlockKernelsTest, GivenKernelWithBlockKernelsWhenProgramIsLink EXPECT_EQ(CL_SUCCESS, retVal); - retVal = programLinked->link(1, &device, buildOptions, 1, &program); + retVal = programLinked->link(pProgram->getDevices(), buildOptions, 1, &program); EXPECT_EQ(CL_SUCCESS, retVal); BlockKernelManager *blockManager = programLinked->getBlockKernelManager(); diff --git a/opencl/test/unit_test/program/program_with_kernel_debug_tests.cpp b/opencl/test/unit_test/program/program_with_kernel_debug_tests.cpp index 7205243c87..ea13a0e9fd 100644 --- a/opencl/test/unit_test/program/program_with_kernel_debug_tests.cpp +++ b/opencl/test/unit_test/program/program_with_kernel_debug_tests.cpp @@ -185,7 +185,7 @@ TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsLinke EXPECT_CALL(*program, appendKernelDebugOptions()).Times(1); cl_program clProgramToLink = pProgram; - retVal = program->link(1, &device, nullptr, 1, &clProgramToLink); + retVal = program->link(pProgram->getDevices(), nullptr, 1, &clProgramToLink); EXPECT_EQ(CL_SUCCESS, retVal); } @@ -222,7 +222,7 @@ TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsLinke EXPECT_EQ(CL_SUCCESS, retVal); cl_program program = pProgram; - retVal = pProgram->link(1, &device, nullptr, + retVal = pProgram->link(pProgram->getDevices(), nullptr, 1, &program); EXPECT_EQ(CL_SUCCESS, retVal); }