Media/compute-runtime
1
0
Fork 0
mirror of https://github.com/intel/compute-runtime.git synced 2025-09-15 13:01:45 +08:00
Code Issues Packages Projects Releases Wiki Activity

Use all passed devices when linking program

Browse Source 
Related-To: NEO-5001
Change-Id: I3944375023b9cb3322df83e6e042ba8be2bd762c
Signed-off-by: Mateusz Jablonski <mateusz.jablonski@intel.com>
This commit is contained in:
Mateusz Jablonski
2020-10-28 17:18:37 +01:00
committed by sys_ocldev
parent 17970ad5e7
commit 430cca3f4c
11 changed files with 397 additions and 119 deletions
Download Patch File Download Diff File Expand all files Collapse all files

47
opencl/source/api/api.cpp
View File

@ -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<ClDevice>(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) {

2
opencl/source/cl_device/cl_device.cpp
View File

@ -169,7 +169,7 @@ ClDeviceVector::ClDeviceVector(const cl_device_id *devices,
}
}
void ClDeviceVector::toDeviceIDs(std::vector<cl_device_id> &devIDs) {
void ClDeviceVector::toDeviceIDs(std::vector<cl_device_id> &devIDs) const {
int i = 0;
devIDs.resize(this->size());

2
opencl/source/cl_device/cl_device_vector.h
View File

@ -20,7 +20,7 @@ class ClDeviceVector : public StackVec<ClDevice *, 1> {
ClDeviceVector &operator=(const ClDeviceVector &) = default;
ClDeviceVector(const cl_device_id *devices,
cl_uint numDevices);
void toDeviceIDs(std::vector<cl_device_id> &devIDs);
void toDeviceIDs(std::vector<cl_device_id> &devIDs) const;
};
} // namespace NEO

100
opencl/source/program/link.cpp
View File

@ -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<ClDevice>();
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<ClDevice>();
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();

24
opencl/source/program/program.cpp
View File

@ -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<ClDeviceVector *>(&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<ClDevice>(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

4
opencl/source/program/program.h
View File

@ -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);

175
opencl/test/unit_test/api/cl_get_program_build_info_tests.inl
View File

@ -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<char[]> 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<MockProgram>(
&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<char[]> 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<MockProgram>(
&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<cl_device_id> 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

108
opencl/test/unit_test/api/cl_link_program_tests.inl
View File

@ -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

48
opencl/test/unit_test/program/program_tests.cpp
View File

@ -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<ClDevice *>(usedDevice);
auto p2 = std::make_unique<FailingGenBinaryProgram>(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<MockProgram>(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<MockProgram>(
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<void *>(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();

2
opencl/test/unit_test/program/program_with_block_kernels_tests.cpp
View File

@ -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();

4
opencl/test/unit_test/program/program_with_kernel_debug_tests.cpp
View File

@ -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);
}