Media/compute-runtime
1
0
Fork 0
mirror of https://github.com/intel/compute-runtime.git synced 2025-09-10 12:53:42 +08:00
Code Issues Packages Projects Releases Wiki Activity

Use all devices when building program

Browse Source 
remove Program::internalOptions
internal options are calculated separately in compile, link and build methods

Related-To: NEO-5001
Change-Id: I85ea2d64c72edb9b0e3f23244b59b9af20e2d357
Signed-off-by: Mateusz Jablonski <mateusz.jablonski@intel.com>
This commit is contained in:
Mateusz Jablonski
2020-10-30 11:10:00 +01:00
committed by sys_ocldev
parent 045632f355
commit 75d1ebb811
34 changed files with 655 additions and 390 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@ -1511,8 +1511,14 @@ cl_int CL_API_CALL clBuildProgram(cl_program program,
retVal = validateObjects(WithCastToInternal(program, &pProgram), Program::isValidCallback(funcNotify, userData));
ClDeviceVector deviceVector;
ClDeviceVector *deviceVectorPtr = &deviceVector;
if (CL_SUCCESS == retVal) {
retVal = pProgram->build(numDevices, deviceList, options, clCacheEnabled);
retVal = Program::processInputDevices(deviceVectorPtr, numDevices, deviceList, pProgram->getDevices());
}
if (CL_SUCCESS == retVal) {
retVal = pProgram->build(*deviceVectorPtr, options, clCacheEnabled);
pProgram->invokeCallback(funcNotify, userData);
}

2
opencl/source/built_ins/populate_built_ins.inl
View File

@ -14,7 +14,7 @@ void BuiltinDispatchInfoBuilder::populate(ClDevice &device, EBuiltInOps::Type op
ClDeviceVector deviceVector;
deviceVector.push_back(&device);
prog.reset(BuiltinDispatchInfoBuilder::createProgramFromCode(src, deviceVector).release());
prog->build(0, nullptr, options.data(), kernelsLib.isCacheingEnabled());
prog->build(deviceVector, options.data(), kernelsLib.isCacheingEnabled());
grabKernels(std::forward<KernelsDescArgsT>(desc)...);
}
} // namespace NEO

2
opencl/source/built_ins/vme_builtin.cpp
View File

@ -81,7 +81,7 @@ Program *Vme::createBuiltInProgram(
builtinsBuilders["block_advanced_motion_estimate_bidirectional_check_intel"] =
&Vme::getBuiltinDispatchInfoBuilder(EBuiltInOps::VmeBlockAdvancedMotionEstimateBidirectionalCheckIntel, device);
errcodeRet = pBuiltInProgram->build(&device.getDevice(), mediaKernelsBuildOptions, true, builtinsBuilders);
errcodeRet = pBuiltInProgram->build(deviceVector, mediaKernelsBuildOptions, true, builtinsBuilders);
} else {
errcodeRet = CL_INVALID_VALUE;
}

2
opencl/source/program/additional_options.cpp
View File

@ -8,6 +8,6 @@
#include "opencl/source/program/program.h"
namespace NEO {
void Program::applyAdditionalOptions() {
void Program::applyAdditionalOptions(std::string &internalOptions) {
}
}; // namespace NEO

182
opencl/source/program/build.cpp
View File

@ -30,45 +30,49 @@
namespace NEO {
cl_int Program::build(
cl_uint numDevices,
const cl_device_id *deviceList,
const ClDeviceVector &deviceVector,
const char *buildOptions,
bool enableCaching) {
cl_int retVal = CL_SUCCESS;
std::string internalOptions;
initInternalOptions(internalOptions);
auto defaultClDevice = deviceVector[0];
UNRECOVERABLE_IF(defaultClDevice == nullptr);
auto &defaultDevice = defaultClDevice->getDevice();
auto clDevice = this->pDevice->getSpecializedDevice<ClDevice>();
UNRECOVERABLE_IF(clDevice == nullptr);
enum class BuildPhase {
Init,
SourceCodeNotification,
BinaryCreation,
BinaryProcessing,
DebugDataNotification
};
std::unordered_map<uint32_t, BuildPhase> phaseReached;
for (const auto &clDevice : deviceVector) {
phaseReached[clDevice->getRootDeviceIndex()] = BuildPhase::Init;
}
do {
if (((deviceList == nullptr) && (numDevices != 0)) ||
((deviceList != nullptr) && (numDevices == 0))) {
retVal = CL_INVALID_VALUE;
break;
}
// if a device_list is specified, make sure it points to our device
// NOTE: a null device_list is ok - it means "all devices"
if (deviceList && validateObject(*deviceList) != CL_SUCCESS) {
retVal = CL_INVALID_DEVICE;
break;
}
// check to see if a previous build request is in progress
if (buildStatuses[clDevice] == CL_BUILD_IN_PROGRESS) {
if (std::any_of(deviceVector.begin(), deviceVector.end(), [&](auto device) { return CL_BUILD_IN_PROGRESS == buildStatuses[device]; })) {
retVal = CL_INVALID_OPERATION;
break;
}
if (isCreatedFromBinary == false) {
buildStatuses[clDevice] = CL_BUILD_IN_PROGRESS;
for (const auto &device : deviceVector) {
buildStatuses[device] = CL_BUILD_IN_PROGRESS;
}
if (nullptr != buildOptions) {
options = buildOptions;
} else if (this->createdFrom != CreatedFrom::BINARY) {
options = "";
}
extractInternalOptions(options);
applyAdditionalOptions();
extractInternalOptions(options, internalOptions);
applyAdditionalOptions(internalOptions);
CompilerInterface *pCompilerInterface = pDevice->getCompilerInterface();
CompilerInterface *pCompilerInterface = defaultDevice.getCompilerInterface();
if (!pCompilerInterface) {
retVal = CL_OUT_OF_HOST_MEMORY;
break;
@ -89,28 +93,29 @@ cl_int Program::build(
if (isKernelDebugEnabled()) {
std::string filename;
appendKernelDebugOptions();
notifyDebuggerWithSourceCode(filename);
if (!filename.empty()) {
// Add "-s" flag first so it will be ignored by clang in case the options already have this flag set.
options = std::string("-s ") + filename + " " + options;
for (const auto &clDevice : deviceVector) {
if (BuildPhase::SourceCodeNotification == phaseReached[clDevice->getRootDeviceIndex()]) {
continue;
}
appendKernelDebugOptions(*clDevice, internalOptions);
notifyDebuggerWithSourceCode(*clDevice, filename);
if (!filename.empty()) {
// Add "-s" flag first so it will be ignored by clang in case the options already have this flag set.
options = std::string("-s ") + filename + " " + options;
}
phaseReached[clDevice->getRootDeviceIndex()] = BuildPhase::SourceCodeNotification;
}
}
if (requiresOpenClCFeatures(options)) {
auto compilerExtensionsWithFeaturesOptions = clDevice->peekCompilerExtensionsWithFeatures();
if (internalOptions.find(compilerExtensionsWithFeaturesOptions) == std::string::npos) {
CompilerOptions::concatenateAppend(internalOptions, compilerExtensionsWithFeaturesOptions);
}
auto compilerFeaturesOptions = clDevice->peekCompilerFeatures();
if (internalOptions.find(compilerFeaturesOptions) == std::string::npos) {
CompilerOptions::concatenateAppend(internalOptions, compilerFeaturesOptions);
}
auto compilerExtensionsWithFeaturesOptions = defaultClDevice->peekCompilerExtensionsWithFeatures();
CompilerOptions::concatenateAppend(internalOptions, compilerExtensionsWithFeaturesOptions);
auto compilerFeaturesOptions = defaultClDevice->peekCompilerFeatures();
CompilerOptions::concatenateAppend(internalOptions, compilerFeaturesOptions);
} else {
auto compilerExtensionsOptions = clDevice->peekCompilerExtensions();
if (internalOptions.find(compilerExtensionsOptions) == std::string::npos) {
CompilerOptions::concatenateAppend(internalOptions, compilerExtensionsOptions);
}
auto compilerExtensionsOptions = defaultClDevice->peekCompilerExtensions();
CompilerOptions::concatenateAppend(internalOptions, compilerExtensionsOptions);
}
inputArgs.apiOptions = ArrayRef<const char>(options.c_str(), options.length());
@ -122,28 +127,48 @@ cl_int Program::build(
"\nBuild Internal Options", inputArgs.internalOptions.begin());
inputArgs.allowCaching = enableCaching;
NEO::TranslationOutput compilerOuput = {};
auto compilerErr = pCompilerInterface->build(*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());
retVal = asClError(compilerErr);
for (const auto &clDevice : deviceVector) {
auto compilerErr = pCompilerInterface->build(clDevice->getDevice(), inputArgs, compilerOuput);
this->updateBuildLog(clDevice->getRootDeviceIndex(), compilerOuput.frontendCompilerLog.c_str(), compilerOuput.frontendCompilerLog.size());
this->updateBuildLog(clDevice->getRootDeviceIndex(), compilerOuput.backendCompilerLog.c_str(), compilerOuput.backendCompilerLog.size());
retVal = asClError(compilerErr);
if (retVal != CL_SUCCESS) {
break;
}
if (inputArgs.srcType == IGC::CodeType::oclC) {
this->irBinary = std::move(compilerOuput.intermediateRepresentation.mem);
this->irBinarySize = compilerOuput.intermediateRepresentation.size;
this->isSpirV = compilerOuput.intermediateCodeType == IGC::CodeType::spirV;
}
if (BuildPhase::BinaryCreation == phaseReached[clDevice->getRootDeviceIndex()]) {
continue;
}
this->replaceDeviceBinary(std::move(compilerOuput.deviceBinary.mem), compilerOuput.deviceBinary.size, clDevice->getRootDeviceIndex());
phaseReached[clDevice->getRootDeviceIndex()] = BuildPhase::BinaryCreation;
}
if (retVal != CL_SUCCESS) {
break;
}
if (inputArgs.srcType == IGC::CodeType::oclC) {
this->irBinary = std::move(compilerOuput.intermediateRepresentation.mem);
this->irBinarySize = compilerOuput.intermediateRepresentation.size;
this->isSpirV = compilerOuput.intermediateCodeType == IGC::CodeType::spirV;
}
this->replaceDeviceBinary(std::move(compilerOuput.deviceBinary.mem), compilerOuput.deviceBinary.size, clDevice->getRootDeviceIndex());
this->debugData = std::move(compilerOuput.debugData.mem);
this->debugDataSize = compilerOuput.debugData.size;
}
updateNonUniformFlag();
if (DebugManager.flags.PrintProgramBinaryProcessingTime.get()) {
retVal = TimeMeasureWrapper::functionExecution(*this, &Program::processGenBinary, pDevice->getRootDeviceIndex());
} else {
retVal = processGenBinary(pDevice->getRootDeviceIndex());
for (auto &clDevice : deviceVector) {
if (BuildPhase::BinaryProcessing == phaseReached[clDevice->getRootDeviceIndex()]) {
continue;
}
if (DebugManager.flags.PrintProgramBinaryProcessingTime.get()) {
retVal = TimeMeasureWrapper::functionExecution(*this, &Program::processGenBinary, clDevice->getRootDeviceIndex());
} else {
retVal = processGenBinary(clDevice->getRootDeviceIndex());
}
if (retVal != CL_SUCCESS) {
break;
}
phaseReached[clDevice->getRootDeviceIndex()] = BuildPhase::BinaryProcessing;
}
if (retVal != CL_SUCCESS) {
@ -152,15 +177,20 @@ cl_int Program::build(
if (isKernelDebugEnabled()) {
processDebugData();
auto clDevice = this->pDevice->getSpecializedDevice<ClDevice>();
UNRECOVERABLE_IF(clDevice == nullptr);
if (clDevice->getSourceLevelDebugger()) {
for (auto kernelInfo : kernelInfoArray) {
clDevice->getSourceLevelDebugger()->notifyKernelDebugData(&kernelInfo->debugData,
kernelInfo->kernelDescriptor.kernelMetadata.kernelName,
kernelInfo->heapInfo.pKernelHeap,
kernelInfo->heapInfo.KernelHeapSize);
for (auto &clDevice : deviceVector) {
if (BuildPhase::DebugDataNotification == phaseReached[clDevice->getRootDeviceIndex()]) {
continue;
}
if (clDevice->getSourceLevelDebugger()) {
for (auto kernelInfo : kernelInfoArray) {
clDevice->getSourceLevelDebugger()->notifyKernelDebugData(&kernelInfo->debugData,
kernelInfo->kernelDescriptor.kernelMetadata.kernelName,
kernelInfo->heapInfo.pKernelHeap,
kernelInfo->heapInfo.KernelHeapSize);
}
}
phaseReached[clDevice->getRootDeviceIndex()] = BuildPhase::DebugDataNotification;
}
}
@ -168,43 +198,39 @@ cl_int Program::build(
} 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;
programBinaryType = CL_PROGRAM_BINARY_TYPE_EXECUTABLE;
for (const auto &device : deviceVector) {
buildStatuses[device] = CL_BUILD_SUCCESS;
}
}
return retVal;
}
bool Program::appendKernelDebugOptions() {
bool Program::appendKernelDebugOptions(ClDevice &clDevice, std::string &internalOptions) {
CompilerOptions::concatenateAppend(internalOptions, CompilerOptions::debugKernelEnable);
CompilerOptions::concatenateAppend(options, CompilerOptions::generateDebugInfo);
auto clDevice = this->pDevice->getSpecializedDevice<ClDevice>();
UNRECOVERABLE_IF(clDevice == nullptr);
auto debugger = clDevice->getSourceLevelDebugger();
auto debugger = clDevice.getSourceLevelDebugger();
if (debugger && debugger->isOptimizationDisabled()) {
CompilerOptions::concatenateAppend(options, CompilerOptions::optDisable);
}
return true;
}
void Program::notifyDebuggerWithSourceCode(std::string &filename) {
auto clDevice = this->pDevice->getSpecializedDevice<ClDevice>();
UNRECOVERABLE_IF(clDevice == nullptr);
if (clDevice->getSourceLevelDebugger()) {
clDevice->getSourceLevelDebugger()->notifySourceCode(sourceCode.c_str(), sourceCode.size(), filename);
void Program::notifyDebuggerWithSourceCode(ClDevice &clDevice, std::string &filename) {
if (clDevice.getSourceLevelDebugger()) {
clDevice.getSourceLevelDebugger()->notifySourceCode(sourceCode.c_str(), sourceCode.size(), filename);
}
}
cl_int Program::build(const Device *pDevice, const char *buildOptions, bool enableCaching,
cl_int Program::build(const ClDeviceVector &deviceVector, const char *buildOptions, bool enableCaching,
std::unordered_map<std::string, BuiltinDispatchInfoBuilder *> &builtinsMap) {
cl_device_id deviceId = pDevice->getSpecializedDevice<ClDevice>();
auto ret = this->build(1, &deviceId, buildOptions, enableCaching);
auto ret = this->build(deviceVector, buildOptions, enableCaching);
if (ret != CL_SUCCESS) {
return ret;
}
@ -219,7 +245,7 @@ cl_int Program::build(const Device *pDevice, const char *buildOptions, bool enab
return ret;
}
void Program::extractInternalOptions(const std::string &options) {
void Program::extractInternalOptions(const std::string &options, std::string &internalOptions) {
auto tokenized = CompilerOptions::tokenize(options);
for (auto &optionString : internalOptionsToExtract) {
auto element = std::find(tokenized.begin(), tokenized.end(), optionString);

23
opencl/source/program/compile.cpp
View File

@ -36,7 +36,9 @@ cl_int Program::compile(
auto defaultClDevice = deviceVector[0];
UNRECOVERABLE_IF(defaultClDevice == nullptr);
auto &defaultDevice = defaultClDevice->getDevice();
internalOptions.clear();
std::string internalOptions;
initInternalOptions(internalOptions);
std::unordered_map<uint32_t, bool> sourceLevelDebuggerNotified;
do {
if (numInputHeaders == 0) {
if ((headerIncludeNames != nullptr) || (inputHeaders != nullptr)) {
@ -60,6 +62,7 @@ cl_int Program::compile(
break;
}
for (const auto &device : deviceVector) {
sourceLevelDebuggerNotified[device->getRootDeviceIndex()] = false;
buildStatuses[device] = CL_BUILD_IN_PROGRESS;
}
@ -121,11 +124,18 @@ cl_int Program::compile(
}
if (isKernelDebugEnabled()) {
std::string filename;
appendKernelDebugOptions();
notifyDebuggerWithSourceCode(filename);
if (!filename.empty()) {
options = std::string("-s ") + filename + " " + options;
for (const auto &device : deviceVector) {
if (sourceLevelDebuggerNotified[device->getRootDeviceIndex()]) {
continue;
}
appendKernelDebugOptions(*device, internalOptions);
std::string filename;
notifyDebuggerWithSourceCode(*device, filename);
if (!filename.empty()) {
options = std::string("-s ") + filename + " " + options;
}
sourceLevelDebuggerNotified[device->getRootDeviceIndex()] = true;
}
}
@ -165,7 +175,6 @@ cl_int Program::compile(
programBinaryType = CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT;
}
internalOptions.clear();
return retVal;
}
} // namespace NEO

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

@ -38,7 +38,10 @@ cl_int Program::link(
auto defaultClDevice = deviceVector[0];
UNRECOVERABLE_IF(defaultClDevice == nullptr);
auto &defaultDevice = defaultClDevice->getDevice();
internalOptions.clear();
std::unordered_map<uint32_t, bool> kernelDebugDataNotified;
std::unordered_map<uint32_t, bool> debugOptionsAppended;
std::string internalOptions;
initInternalOptions(internalOptions);
do {
if ((numInputPrograms == 0) || (inputPrograms == nullptr)) {
retVal = CL_INVALID_VALUE;
@ -51,6 +54,8 @@ cl_int Program::link(
}
for (const auto &device : deviceVector) {
kernelDebugDataNotified[device->getRootDeviceIndex()] = false;
debugOptionsAppended[device->getRootDeviceIndex()] = false;
buildStatuses[device] = CL_BUILD_IN_PROGRESS;
}
@ -65,7 +70,14 @@ cl_int Program::link(
}
if (isKernelDebugEnabled()) {
appendKernelDebugOptions();
for (auto &device : deviceVector) {
if (debugOptionsAppended[device->getRootDeviceIndex()]) {
continue;
}
appendKernelDebugOptions(*device, internalOptions);
debugOptionsAppended[device->getRootDeviceIndex()] = true;
}
}
isCreateLibrary = CompilerOptions::contains(options, CompilerOptions::createLibrary);
@ -154,6 +166,9 @@ cl_int Program::link(
programBinaryType = CL_PROGRAM_BINARY_TYPE_EXECUTABLE;
if (isKernelDebugEnabled()) {
if (kernelDebugDataNotified[device->getRootDeviceIndex()]) {
continue;
}
processDebugData();
for (auto kernelInfo : kernelInfoArray) {
device->getSourceLevelDebugger()->notifyKernelDebugData(&kernelInfo->debugData,
@ -161,6 +176,7 @@ cl_int Program::link(
kernelInfo->heapInfo.pKernelHeap,
kernelInfo->heapInfo.KernelHeapSize);
}
kernelDebugDataNotified[device->getRootDeviceIndex()] = true;
}
}
@ -201,8 +217,6 @@ cl_int Program::link(
}
}
internalOptions.clear();
return retVal;
}
} // namespace NEO

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

@ -48,7 +48,6 @@ Program::Program(Context *context, bool isBuiltIn, const ClDeviceVector &clDevic
ClDevice *pClDevice = castToObject<ClDevice>(pDevice->getSpecializedDevice<ClDevice>());
numDevices = static_cast<uint32_t>(clDevicesIn.size());
bool force32BitAddressess = false;
uint32_t maxRootDeviceIndex = 0;
@ -60,7 +59,11 @@ Program::Program(Context *context, bool isBuiltIn, const ClDeviceVector &clDevic
}
buildInfos.resize(maxRootDeviceIndex + 1);
kernelDebugEnabled = pClDevice->isDebuggerActive();
}
void Program::initInternalOptions(std::string &internalOptions) const {
auto pClDevice = clDevices[0];
auto force32BitAddressess = pClDevice->getSharedDeviceInfo().force32BitAddressess;
auto enabledClVersion = pClDevice->getEnabledClVersion();
if (enabledClVersion == 30) {
internalOptions = "-ocl-version=300 ";
@ -69,7 +72,6 @@ Program::Program(Context *context, bool isBuiltIn, const ClDeviceVector &clDevic
} else {
internalOptions = "-ocl-version=120 ";
}
force32BitAddressess = pClDevice->getSharedDeviceInfo().force32BitAddressess;
if (force32BitAddressess && !isBuiltIn) {
CompilerOptions::concatenateAppend(internalOptions, CompilerOptions::arch32bit);
@ -85,8 +87,6 @@ Program::Program(Context *context, bool isBuiltIn, const ClDeviceVector &clDevic
CompilerOptions::concatenateAppend(internalOptions, CompilerOptions::bindlessImages);
}
kernelDebugEnabled = pClDevice->isDebuggerActive();
auto enableStatelessToStatefullWithOffset = pClDevice->getHardwareCapabilities().isStatelesToStatefullWithOffsetSupported;
if (DebugManager.flags.EnableStatelessToStatefulBufferOffsetOpt.get() != -1) {
enableStatelessToStatefullWithOffset = DebugManager.flags.EnableStatelessToStatefulBufferOffsetOpt.get() != 0;

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

@ -126,10 +126,10 @@ class Program : public BaseObject<_cl_program> {
Program(const Program &) = delete;
Program &operator=(const Program &) = delete;
cl_int build(cl_uint numDevices, const cl_device_id *deviceList, const char *buildOptions,
cl_int build(const ClDeviceVector &deviceVector, const char *buildOptions,
bool enableCaching);
cl_int build(const Device *pDevice, const char *buildOptions, bool enableCaching,
cl_int build(const ClDeviceVector &deviceVector, const char *buildOptions, bool enableCaching,
std::unordered_map<std::string, BuiltinDispatchInfoBuilder *> &builtinsMap);
MOCKABLE_VIRTUAL cl_int processGenBinary(uint32_t rootDeviceIndex);
@ -215,8 +215,6 @@ class Program : public BaseObject<_cl_program> {
const std::string &getOptions() const { return options; }
const std::string &getInternalOptions() const { return internalOptions; }
bool getAllowNonUniform() const {
return allowNonUniform;
}
@ -267,6 +265,7 @@ class Program : public BaseObject<_cl_program> {
bool isDeviceAssociated(const ClDevice &clDevice) const;
static cl_int processInputDevices(ClDeviceVector *&deviceVectorPtr, cl_uint numDevices, const cl_device_id *deviceList, const ClDeviceVector &allAvailableDevices);
MOCKABLE_VIRTUAL void initInternalOptions(std::string &internalOptions) const;
protected:
MOCKABLE_VIRTUAL cl_int createProgramFromBinary(const void *pBinary, size_t binarySize, uint32_t rootDeviceIndex);
@ -280,13 +279,13 @@ class Program : public BaseObject<_cl_program> {
void updateNonUniformFlag();
void updateNonUniformFlag(const Program **inputProgram, size_t numInputPrograms);
void extractInternalOptions(const std::string &options);
void extractInternalOptions(const std::string &options, std::string &internalOptions);
MOCKABLE_VIRTUAL bool isFlagOption(ConstStringRef option);
MOCKABLE_VIRTUAL bool isOptionValueValid(ConstStringRef option, ConstStringRef value);
MOCKABLE_VIRTUAL void applyAdditionalOptions();
MOCKABLE_VIRTUAL void applyAdditionalOptions(std::string &internalOptions);
MOCKABLE_VIRTUAL bool appendKernelDebugOptions();
void notifyDebuggerWithSourceCode(std::string &filename);
MOCKABLE_VIRTUAL bool appendKernelDebugOptions(ClDevice &clDevice, std::string &internalOptions);
void notifyDebuggerWithSourceCode(ClDevice &clDevice, std::string &filename);
void setBuildStatus(cl_build_status status);
@ -310,7 +309,6 @@ class Program : public BaseObject<_cl_program> {
std::string sourceCode;
std::string options;
std::string internalOptions;
static const std::vector<ConstStringRef> internalOptionsToExtract;
uint32_t programOptionVersion = 12U;

74
opencl/test/unit_test/api/cl_build_program_tests.inl
View File

@ -309,4 +309,78 @@ TEST_F(clBuildProgramTests, GivenValidCallbackInputWhenBuildProgramThenCallbackI
EXPECT_EQ(CL_SUCCESS, retVal);
}
TEST_F(clBuildProgramTests, givenMultiDeviceProgramWhenBuildingForInvalidDevicesInputThenInvalidDeviceErrorIsReturned) {
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);
MockContext mockContext;
cl_device_id nullDeviceInput[] = {pContext->getDevice(0), nullptr};
cl_device_id notAssociatedDeviceInput[] = {mockContext.getDevice(0)};
cl_device_id validDeviceInput[] = {pContext->getDevice(0)};
retVal = clBuildProgram(
pProgram,
0,
validDeviceInput,
nullptr,
nullptr,
nullptr);
EXPECT_EQ(CL_INVALID_VALUE, retVal);
retVal = clBuildProgram(
pProgram,
1,
nullptr,
nullptr,
nullptr,
nullptr);
EXPECT_EQ(CL_INVALID_VALUE, retVal);
retVal = clBuildProgram(
pProgram,
2,
nullDeviceInput,
nullptr,
nullptr,
nullptr);
EXPECT_EQ(CL_INVALID_DEVICE, retVal);
retVal = clBuildProgram(
pProgram,
1,
notAssociatedDeviceInput,
nullptr,
nullptr,
nullptr);
EXPECT_EQ(CL_INVALID_DEVICE, retVal);
retVal = clReleaseProgram(pProgram);
EXPECT_EQ(CL_SUCCESS, retVal);
}
} // namespace ULT

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

@ -451,4 +451,133 @@ TEST(clGetProgramBuildInfoTest, givenMultiDeviceProgramWhenLinkingWithoutInputDe
retVal = clReleaseProgram(outProgram);
EXPECT_EQ(CL_SUCCESS, retVal);
}
TEST(clGetProgramBuildInfoTest, givenMultiDeviceProgramWhenBuildingForSpecificDevicesThenOnlySpecificDevicesReportBuildStatus) {
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);
}
cl_device_id devicesForBuild[] = {context.getDevice(1), context.getDevice(3)};
cl_device_id devicesNotForBuild[] = {context.getDevice(0), context.getDevice(2), context.getDevice(4), context.getDevice(5)};
retVal = clBuildProgram(
pProgram,
2,
devicesForBuild,
nullptr,
nullptr,
nullptr);
ASSERT_EQ(CL_SUCCESS, retVal);
for (const auto &device : devicesForBuild) {
retVal = clGetProgramBuildInfo(pProgram, device, CL_PROGRAM_BUILD_STATUS, sizeof(buildStatus), &buildStatus, NULL);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_EQ(CL_BUILD_SUCCESS, buildStatus);
}
for (const auto &device : devicesNotForBuild) {
retVal = clGetProgramBuildInfo(pProgram, device, CL_PROGRAM_BUILD_STATUS, sizeof(buildStatus), &buildStatus, NULL);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_EQ(CL_BUILD_NONE, buildStatus);
}
retVal = clReleaseProgram(pProgram);
EXPECT_EQ(CL_SUCCESS, retVal);
}
TEST(clGetProgramBuildInfoTest, givenMultiDeviceProgramWhenBuildingWithoutInputDevicesThenAllDevicesReportBuildStatus) {
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 = clBuildProgram(
pProgram,
0,
nullptr,
nullptr,
nullptr,
nullptr);
ASSERT_EQ(CL_SUCCESS, retVal);
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_SUCCESS, buildStatus);
}
retVal = clReleaseProgram(pProgram);
EXPECT_EQ(CL_SUCCESS, retVal);
}
} // namespace ULT

7
opencl/test/unit_test/aub_tests/command_queue/enqueue_kernel_aub_tests.cpp
View File

@ -502,13 +502,6 @@ HWTEST_F(AUBSimpleKernelStatelessTest, givenSimpleKernelWhenStatelessPathIsUsedT
event);
ASSERT_EQ(CL_SUCCESS, retVal);
EXPECT_THAT(this->pProgram->getInternalOptions(),
testing::HasSubstr(std::string(NEO::CompilerOptions::greaterThan4gbBuffersRequired)));
if (this->device->getSharedDeviceInfo().force32BitAddressess) {
EXPECT_THAT(this->pProgram->getInternalOptions(),
testing::HasSubstr(std::string(NEO::CompilerOptions::arch32bit)));
}
EXPECT_FALSE(this->kernel->getKernelInfo().kernelArgInfo[0].pureStatefulBufferAccess);
EXPECT_TRUE(this->kernel->getKernelInfo().patchInfo.executionEnvironment->CompiledForGreaterThan4GBBuffers);

9
opencl/test/unit_test/aub_tests/gen9/skl/command_queue/run_kernel_aub_tests_skl.cpp
View File

@ -42,10 +42,8 @@ SKLTEST_F(AUBRunKernelIntegrateTest, ooqExecution) {
Program *pProgram = CreateProgramFromBinary(kernelFilename);
ASSERT_NE(nullptr, pProgram);
cl_device_id device = pClDevice;
retVal = pProgram->build(
1,
&device,
pProgram->getDevices(),
nullptr,
false);
ASSERT_EQ(CL_SUCCESS, retVal);
@ -267,11 +265,8 @@ SKLTEST_F(AUBRunKernelIntegrateTest, deviceSideVme) {
Program *pProgram = CreateProgramFromBinary(kernelFilename);
ASSERT_NE(nullptr, pProgram);
cl_device_id device = pClDevice;
retVal = pProgram->build(
1,
&device,
pProgram->getDevices(),
"",
false);
ASSERT_EQ(CL_SUCCESS, retVal);

7
opencl/test/unit_test/built_ins/built_in_tests.cpp
View File

@ -1381,7 +1381,9 @@ TEST_F(BuiltInTests, givenCreateProgramFromSourceWhenDeviceSupportSharedSystemAl
EXPECT_NE(0u, bc.resource.size());
auto program = std::unique_ptr<Program>(BuiltinDispatchInfoBuilder::createProgramFromCode(bc, toClDeviceVector(*pClDevice)));
EXPECT_NE(nullptr, program.get());
EXPECT_THAT(program->getInternalOptions(), testing::HasSubstr(std::string(CompilerOptions::greaterThan4gbBuffersRequired)));
std::string builtinInternalOptions;
program->initInternalOptions(builtinInternalOptions);
EXPECT_THAT(builtinInternalOptions, testing::HasSubstr(std::string(CompilerOptions::greaterThan4gbBuffersRequired)));
}
TEST_F(BuiltInTests, GivenTypeIntermediateWhenCreatingProgramFromCodeThenNullPointerIsReturned) {
@ -1426,7 +1428,8 @@ TEST_F(BuiltInTests, GivenForce32bitWhenCreatingProgramThenCorrectKernelIsCreate
auto program = std::unique_ptr<Program>(BuiltinDispatchInfoBuilder::createProgramFromCode(bc, toClDeviceVector(*pClDevice)));
ASSERT_NE(nullptr, program.get());
auto builtinInternalOptions = program->getInternalOptions();
std::string builtinInternalOptions;
program->initInternalOptions(builtinInternalOptions);
auto it = builtinInternalOptions.find(NEO::CompilerOptions::arch32bit.data());
EXPECT_EQ(std::string::npos, it);

3
opencl/test/unit_test/command_queue/enqueue_debug_kernel_tests.cpp
View File

@ -44,11 +44,10 @@ class EnqueueDebugKernelTest : public ProgramSimpleFixture,
kbHelper = new KernelBinaryHelper(filename, false);
CreateProgramWithSource(
pContext,
&device,
"copybuffer.cl");
pProgram->enableKernelDebug();
cl_int retVal = pProgram->build(1, &device, nullptr, false);
cl_int retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
ASSERT_EQ(CL_SUCCESS, retVal);
// create a kernel

6
opencl/test/unit_test/command_queue/enqueue_svm_tests.cpp
View File

@ -737,8 +737,7 @@ TEST_F(EnqueueSvmTest, GivenSvmAllocationWhenEnqueingKernelThenSuccessIsReturned
EXPECT_NE(nullptr, ptrSVM);
std::unique_ptr<Program> program(Program::createBuiltInFromSource("FillBufferBytes", context, context->getDevices(), &retVal));
cl_device_id device = pClDevice;
program->build(1, &device, nullptr, false);
program->build(program->getDevices(), nullptr, false);
std::unique_ptr<MockKernel> kernel(Kernel::create<MockKernel>(program.get(), *program->getKernelInfo("FillBufferBytes"), &retVal));
kernel->setSvmKernelExecInfo(pSvmAlloc);
@ -766,8 +765,7 @@ TEST_F(EnqueueSvmTest, givenEnqueueTaskBlockedOnUserEventWhenItIsEnqueuedThenSur
EXPECT_NE(nullptr, ptrSVM);
auto program = clUniquePtr(Program::createBuiltInFromSource("FillBufferBytes", context, context->getDevices(), &retVal));
cl_device_id device = pClDevice;
program->build(1, &device, nullptr, false);
program->build(program->getDevices(), nullptr, false);
auto kernel = clUniquePtr(Kernel::create<MockKernel>(program.get(), *program->getKernelInfo("FillBufferBytes"), &retVal));
std::vector<Surface *> allSurfaces;

3
opencl/test/unit_test/command_queue/get_size_required_image_tests.cpp
View File

@ -116,8 +116,7 @@ HWTEST_F(GetSizeRequiredImageTest, WhenCopyingReadWriteImageThenHeapsAndCommandB
auto usedBeforeSSH = ssh.getUsed();
std::unique_ptr<Program> program(Program::createBuiltInFromSource("CopyImageToImage3d", context, context->getDevices(), nullptr));
cl_device_id device = pClDevice;
program->build(1, &device, nullptr, false);
program->build(program->getDevices(), nullptr, false);
std::unique_ptr<Kernel> kernel(Kernel::create<MockKernel>(program.get(), *program->getKernelInfo("CopyImageToImage3d"), nullptr));
EXPECT_NE(nullptr, kernel);

6
opencl/test/unit_test/context/driver_diagnostics_tests.h
View File

@ -222,8 +222,7 @@ struct PerformanceHintEnqueueKernelTest : public PerformanceHintEnqueueTest,
void SetUp() override {
PerformanceHintEnqueueTest::SetUp();
CreateProgramFromBinary(context, context->getDevices(), "CopyBuffer_simd32");
cl_device_id device = context->getDevice(0);
retVal = pProgram->build(1, &device, nullptr, false);
retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
ASSERT_EQ(CL_SUCCESS, retVal);
kernel = Kernel::create<MockKernel>(pProgram, *pProgram->getKernelInfo("CopyBuffer"), &retVal);
@ -257,9 +256,8 @@ struct PerformanceHintEnqueueKernelPrintfTest : public PerformanceHintEnqueueTes
void SetUp() override {
PerformanceHintEnqueueTest::SetUp();
cl_device_id device = pPlatform->getClDevice(0);
CreateProgramFromBinary(context, context->getDevices(), "printf");
retVal = pProgram->build(1, &device, nullptr, false);
retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
ASSERT_EQ(CL_SUCCESS, retVal);
kernel = Kernel::create(pProgram, *pProgram->getKernelInfo("test"), &retVal);

4
opencl/test/unit_test/fixtures/execution_model_kernel_fixture.h
View File

@ -31,10 +31,8 @@ class ExecutionModelKernelFixture : public ProgramFromBinaryTest,
ASSERT_NE(nullptr, pProgram);
ASSERT_EQ(CL_SUCCESS, retVal);
cl_device_id device = pClDevice;
retVal = pProgram->build(
1,
&device,
pProgram->getDevices(),
nullptr,
false);
ASSERT_EQ(CL_SUCCESS, retVal);

5
opencl/test/unit_test/fixtures/hello_world_kernel_fixture.h
View File

@ -68,11 +68,8 @@ struct HelloWorldKernelFixture : public ProgramFixture {
ASSERT_NE(nullptr, pProgram);
cl_device_id device = pDevice;
retVal = pProgram->build(
1,
&device,
pProgram->getDevices(),
nullptr,
false);
ASSERT_EQ(CL_SUCCESS, retVal);

1
opencl/test/unit_test/fixtures/program_fixture.cpp
View File

@ -12,7 +12,6 @@
namespace NEO {
void ProgramFixture::CreateProgramWithSource(Context *pContext,
cl_device_id *deviceList,
const std::string &sourceFileName) {
Cleanup();
cl_int retVal = CL_SUCCESS;

1
opencl/test/unit_test/fixtures/program_fixture.h
View File

@ -30,7 +30,6 @@ class ProgramFixture {
const std::string &options = "");
void CreateProgramWithSource(Context *pContext,
cl_device_id *pDeviceList,
const std::string &sourceFileName);
protected:

21
opencl/test/unit_test/fixtures/simple_arg_kernel_fixture.h
View File

@ -90,7 +90,6 @@ class SimpleArgKernelFixture : public ProgramFixture {
testFile.replace(pos, 1, "_");
}
cl_device_id device = pDevice;
auto deviceVector = toClDeviceVector(*pDevice);
pContext = Context::create<MockContext>(nullptr, deviceVector, nullptr, nullptr, retVal);
ASSERT_EQ(CL_SUCCESS, retVal);
@ -103,8 +102,7 @@ class SimpleArgKernelFixture : public ProgramFixture {
ASSERT_NE(nullptr, pProgram);
retVal = pProgram->build(
1,
&device,
pProgram->getDevices(),
nullptr,
false);
ASSERT_EQ(CL_SUCCESS, retVal);
@ -143,8 +141,6 @@ class SimpleArgNonUniformKernelFixture : public ProgramFixture {
void SetUp(ClDevice *device, Context *context) {
ProgramFixture::SetUp();
cl_device_id deviceId = device;
CreateProgramFromBinary(
context,
context->getDevices(),
@ -153,8 +149,7 @@ class SimpleArgNonUniformKernelFixture : public ProgramFixture {
ASSERT_NE(nullptr, pProgram);
retVal = pProgram->build(
1,
&deviceId,
pProgram->getDevices(),
"-cl-std=CL2.0",
false);
ASSERT_EQ(CL_SUCCESS, retVal);
@ -188,7 +183,6 @@ class SimpleKernelFixture : public ProgramFixture {
void SetUp(ClDevice *device, Context *context) {
ProgramFixture::SetUp();
cl_device_id deviceId = device;
std::string programName("simple_kernels");
CreateProgramFromBinary(
context,
@ -197,8 +191,7 @@ class SimpleKernelFixture : public ProgramFixture {
ASSERT_NE(nullptr, pProgram);
retVal = pProgram->build(
1,
&deviceId,
pProgram->getDevices(),
nullptr,
false);
ASSERT_EQ(CL_SUCCESS, retVal);
@ -241,7 +234,6 @@ class SimpleKernelStatelessFixture : public ProgramFixture {
protected:
void SetUp(ClDevice *device, Context *context) {
ProgramFixture::SetUp();
cl_device_id deviceId = device;
DebugManager.flags.DisableStatelessToStatefulOptimization.set(true);
DebugManager.flags.EnableStatelessToStatefulBufferOffsetOpt.set(false);
@ -252,8 +244,7 @@ class SimpleKernelStatelessFixture : public ProgramFixture {
ASSERT_NE(nullptr, pProgram);
retVal = pProgram->build(
1,
&deviceId,
pProgram->getDevices(),
CompilerOptions::greaterThan4gbBuffersRequired.data(),
false);
ASSERT_EQ(CL_SUCCESS, retVal);
@ -289,7 +280,6 @@ class BindlessKernelFixture : public ProgramFixture {
void createKernel(const std::string &programName, const std::string &kernelName) {
DebugManager.flags.UseBindlessMode.set(1);
cl_device_id deviceId = deviceCl;
CreateProgramFromBinary(
contextCl,
contextCl->getDevices(),
@ -297,8 +287,7 @@ class BindlessKernelFixture : public ProgramFixture {
ASSERT_NE(nullptr, pProgram);
retVal = pProgram->build(
1,
&deviceId,
pProgram->getDevices(),
nullptr,
false);
ASSERT_EQ(CL_SUCCESS, retVal);

4
opencl/test/unit_test/kernel/kernel_arg_info_tests.cpp
View File

@ -31,10 +31,8 @@ class KernelArgInfoTest : public ProgramFromSourceTest {
ASSERT_NE(nullptr, pProgram);
ASSERT_EQ(CL_SUCCESS, retVal);
cl_device_id device = pPlatform->getClDevice(0);
retVal = pProgram->build(
1,
&device,
pProgram->getDevices(),
nullptr,
false);
ASSERT_EQ(CL_SUCCESS, retVal);

3
opencl/test/unit_test/kernel/kernel_is_patched_tests.cpp
View File

@ -22,8 +22,7 @@ class PatchedKernelTest : public ::testing::Test {
context.reset(new MockContext(device.get()));
program.reset(Program::createBuiltInFromSource("FillBufferBytes", context.get(), context->getDevices(), &retVal));
EXPECT_EQ(CL_SUCCESS, retVal);
cl_device_id clDevice = device.get();
program->build(1, &clDevice, nullptr, false);
program->build(program->getDevices(), nullptr, false);
kernel.reset(Kernel::create(program.get(), *program->getKernelInfo("FillBufferBytes"), &retVal));
EXPECT_EQ(CL_SUCCESS, retVal);
}

19
opencl/test/unit_test/kernel/kernel_tests.cpp
View File

@ -58,10 +58,8 @@ class KernelTest : public ProgramFromBinaryTest {
ASSERT_NE(nullptr, pProgram);
ASSERT_EQ(CL_SUCCESS, retVal);
cl_device_id device = pClDevice;
retVal = pProgram->build(
1,
&device,
pProgram->getDevices(),
nullptr,
false);
ASSERT_EQ(CL_SUCCESS, retVal);
@ -343,14 +341,11 @@ class KernelFromBinaryTest : public ProgramSimpleFixture {
typedef Test<KernelFromBinaryTest> KernelFromBinaryTests;
TEST_F(KernelFromBinaryTests, GivenKernelNumArgsWhenGettingInfoThenNumberOfKernelArgsIsReturned) {
cl_device_id device = pClDevice;
CreateProgramFromBinary(pContext, pContext->getDevices(), "kernel_num_args");
ASSERT_NE(nullptr, pProgram);
retVal = pProgram->build(
1,
&device,
pProgram->getDevices(),
nullptr,
false);
@ -384,14 +379,11 @@ TEST_F(KernelFromBinaryTests, GivenKernelNumArgsWhenGettingInfoThenNumberOfKerne
}
TEST_F(KernelFromBinaryTests, WhenRegularKernelIsCreatedThenItIsNotBuiltIn) {
cl_device_id device = pClDevice;
CreateProgramFromBinary(pContext, pContext->getDevices(), "simple_kernels");
ASSERT_NE(nullptr, pProgram);
retVal = pProgram->build(
1,
&device,
pProgram->getDevices(),
nullptr,
false);
@ -417,14 +409,11 @@ TEST_F(KernelFromBinaryTests, WhenRegularKernelIsCreatedThenItIsNotBuiltIn) {
}
TEST_F(KernelFromBinaryTests, givenArgumentDeclaredAsConstantWhenKernelIsCreatedThenArgumentIsMarkedAsReadOnly) {
cl_device_id device = pClDevice;
CreateProgramFromBinary(pContext, pContext->getDevices(), "simple_kernels");
ASSERT_NE(nullptr, pProgram);
retVal = pProgram->build(
1,
&device,
pProgram->getDevices(),
nullptr,
false);

4
opencl/test/unit_test/mocks/mock_program.cpp
View File

@ -30,12 +30,14 @@ ClDeviceVector toClDeviceVector(ClDevice &clDevice) {
}
ProgramInfo *GlobalMockSipProgram::globalSipProgramInfo;
Device *MockProgram::getDevicePtr() { return this->pDevice; }
int MockProgram::initInternalOptionsCalled = 0;
std::string MockProgram::getCachedFileName() const {
auto hwInfo = this->context->getDevice(0)->getHardwareInfo();
auto input = ArrayRef<const char>(this->sourceCode.c_str(), this->sourceCode.size());
auto opts = ArrayRef<const char>(this->options.c_str(), this->options.size());
auto internalOpts = ArrayRef<const char>(this->internalOptions.c_str(), this->internalOptions.size());
auto internalOptions = getInitInternalOptions();
auto internalOpts = ArrayRef<const char>(internalOptions.c_str(), internalOptions.size());
return CompilerCache::getCachedFileName(hwInfo, input, opts, internalOpts);
}

36
opencl/test/unit_test/mocks/mock_program.h
View File

@ -69,7 +69,11 @@ class MockProgram : public Program {
void setBuildOptions(const char *buildOptions) {
options = buildOptions != nullptr ? buildOptions : "";
}
std::string &getInternalOptions() { return internalOptions; };
std::string getInitInternalOptions() const {
std::string internalOptions;
initInternalOptions(internalOptions);
return internalOptions;
};
void setConstantSurface(GraphicsAllocation *gfxAllocation) {
if (gfxAllocation) {
buildInfos[gfxAllocation->getRootDeviceIndex()].constantSurface = gfxAllocation;
@ -143,10 +147,34 @@ class MockProgram : public Program {
this->isCreatedFromBinary = false;
setBuildStatus(CL_BUILD_NONE);
std::unordered_map<std::string, BuiltinDispatchInfoBuilder *> builtins;
auto &device = this->getDevice();
return this->build(&device, this->options.c_str(), false, builtins);
return this->build(getDevices(), this->options.c_str(), false, builtins);
}
void replaceDeviceBinary(std::unique_ptr<char[]> newBinary, size_t newBinarySize, uint32_t rootDeviceIndex) override {
if (replaceDeviceBinaryCalledPerRootDevice.find(rootDeviceIndex) == replaceDeviceBinaryCalledPerRootDevice.end()) {
replaceDeviceBinaryCalledPerRootDevice.insert({rootDeviceIndex, 1});
} else {
replaceDeviceBinaryCalledPerRootDevice[rootDeviceIndex]++;
}
Program::replaceDeviceBinary(std::move(newBinary), newBinarySize, rootDeviceIndex);
}
cl_int processGenBinary(uint32_t rootDeviceIndex) override {
if (processGenBinaryCalledPerRootDevice.find(rootDeviceIndex) == processGenBinaryCalledPerRootDevice.end()) {
processGenBinaryCalledPerRootDevice.insert({rootDeviceIndex, 1});
} else {
processGenBinaryCalledPerRootDevice[rootDeviceIndex]++;
}
return Program::processGenBinary(rootDeviceIndex);
}
void initInternalOptions(std::string &internalOptions) const override {
initInternalOptionsCalled++;
Program::initInternalOptions(internalOptions);
};
std::map<uint32_t, int> processGenBinaryCalledPerRootDevice;
std::map<uint32_t, int> replaceDeviceBinaryCalledPerRootDevice;
static int initInternalOptionsCalled;
bool contextSet = false;
int isFlagOptionOverride = -1;
int isOptionValueValidOverride = -1;
@ -167,7 +195,7 @@ ProgramInfo getSipProgramInfo();
class GMockProgram : public Program {
public:
using Program::Program;
MOCK_METHOD0(appendKernelDebugOptions, bool(void));
MOCK_METHOD(bool, appendKernelDebugOptions, (ClDevice &, std::string &), (override));
};
} // namespace NEO

7
opencl/test/unit_test/program/program_data_tests.cpp
View File

@ -56,7 +56,6 @@ class ProgramDataTestBase : public testing::Test,
CreateProgramWithSource(
pContext,
&device,
"CopyBuffer_simd16.cl");
}
@ -446,8 +445,7 @@ TEST(ProgramScopeMetadataTest, WhenPatchingGlobalSurfaceThenPickProperSourceBuff
}
TEST_F(ProgramDataTest, GivenProgramWith32bitPointerOptWhenProgramScopeConstantBufferPatchTokensAreReadThenConstantPointerOffsetIsPatchedWith32bitPointer) {
cl_device_id device = pPlatform->getClDevice(0);
CreateProgramWithSource(pContext, &device, "CopyBuffer_simd16.cl");
CreateProgramWithSource(pContext, "CopyBuffer_simd16.cl");
ASSERT_NE(nullptr, pProgram);
MockProgram *prog = pProgram;
@ -485,8 +483,7 @@ TEST_F(ProgramDataTest, GivenProgramWith32bitPointerOptWhenProgramScopeConstantB
}
TEST_F(ProgramDataTest, GivenProgramWith32bitPointerOptWhenProgramScopeGlobalPointerPatchTokensAreReadThenGlobalPointerOffsetIsPatchedWith32bitPointer) {
cl_device_id device = pPlatform->getClDevice(0);
CreateProgramWithSource(pContext, &device, "CopyBuffer_simd16.cl");
CreateProgramWithSource(pContext, "CopyBuffer_simd16.cl");
ASSERT_NE(nullptr, pProgram);
MockProgram *prog = pProgram;

9
opencl/test/unit_test/program/program_nonuniform.cpp
View File

@ -190,8 +190,7 @@ TEST_F(ProgramNonUniformTest, GivenCl21WhenExecutingKernelWithNonUniformThenEnqu
mockProgram->setBuildOptions("-cl-std=CL2.1");
retVal = mockProgram->build(
1,
&device,
mockProgram->getDevices(),
nullptr,
false);
EXPECT_EQ(CL_SUCCESS, retVal);
@ -230,8 +229,7 @@ TEST_F(ProgramNonUniformTest, GivenCl20WhenExecutingKernelWithNonUniformThenEnqu
mockProgram->setBuildOptions("-cl-std=CL2.0");
retVal = mockProgram->build(
1,
&device,
mockProgram->getDevices(),
nullptr,
false);
EXPECT_EQ(CL_SUCCESS, retVal);
@ -268,8 +266,7 @@ TEST_F(ProgramNonUniformTest, GivenCl12WhenExecutingKernelWithNonUniformThenInva
mockProgram->setBuildOptions("-cl-std=CL1.2");
retVal = mockProgram->build(
1,
&device,
mockProgram->getDevices(),
nullptr,
false);
EXPECT_EQ(CL_SUCCESS, retVal);

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

@ -109,10 +109,8 @@ class SucceedingGenBinaryProgram : public MockProgram {
};
TEST_P(ProgramFromBinaryTest, WhenBuildingProgramThenSuccessIsReturned) {
cl_device_id device = pClDevice;
retVal = pProgram->build(
1,
&device,
pProgram->getDevices(),
nullptr,
false);
@ -220,11 +218,9 @@ TEST_P(ProgramFromBinaryTest, GivenProgramWithOneKernelWhenGettingNumKernelsThen
size_t paramValue = 0;
size_t paramValueSize = sizeof(paramValue);
size_t paramValueSizeRet = 0;
cl_device_id device = pClDevice;
retVal = pProgram->build(
1,
&device,
pProgram->getDevices(),
nullptr,
false);
ASSERT_EQ(CL_SUCCESS, retVal);
@ -260,11 +256,9 @@ TEST_P(ProgramFromBinaryTest, GivenProgramWithNoExecutableCodeWhenGettingNumKern
TEST_P(ProgramFromBinaryTest, WhenGettingKernelNamesThenCorrectNameIsReturned) {
size_t paramValueSize = sizeof(size_t *);
size_t paramValueSizeRet = 0;
cl_device_id device = pClDevice;
retVal = pProgram->build(
1,
&device,
pProgram->getDevices(),
nullptr,
false);
ASSERT_EQ(CL_SUCCESS, retVal);
@ -593,8 +587,7 @@ TEST_P(ProgramFromBinaryTest, GivenGlobalVariableTotalSizeSetWhenGettingBuildGlo
}
TEST_P(ProgramFromBinaryTest, givenProgramWhenItIsBeingBuildThenItContainsGraphicsAllocationInKernelInfo) {
cl_device_id device = pClDevice;
pProgram->build(1, &device, nullptr, true);
pProgram->build(pProgram->getDevices(), nullptr, true);
auto kernelInfo = pProgram->getKernelInfo(size_t(0));
auto graphicsAllocation = kernelInfo->getGraphicsAllocation();
@ -610,8 +603,7 @@ TEST_P(ProgramFromBinaryTest, givenProgramWhenItIsBeingBuildThenItContainsGraphi
}
TEST_P(ProgramFromBinaryTest, whenProgramIsBeingRebuildThenOutdatedGlobalBuffersAreFreed) {
cl_device_id device = pClDevice;
pProgram->build(1, &device, nullptr, true);
pProgram->build(pProgram->getDevices(), nullptr, true);
EXPECT_EQ(nullptr, pProgram->buildInfos[pClDevice->getRootDeviceIndex()].constantSurface);
EXPECT_EQ(nullptr, pProgram->buildInfos[pClDevice->getRootDeviceIndex()].globalSurface);
@ -627,18 +619,16 @@ TEST_P(ProgramFromBinaryTest, whenProgramIsBeingRebuildThenOutdatedGlobalBuffers
}
TEST_P(ProgramFromBinaryTest, givenProgramWhenCleanKernelInfoIsCalledThenKernelAllocationIsFreed) {
cl_device_id device = pClDevice;
pProgram->build(1, &device, nullptr, true);
pProgram->build(pProgram->getDevices(), nullptr, true);
EXPECT_EQ(1u, pProgram->getNumKernels());
pProgram->cleanCurrentKernelInfo();
EXPECT_EQ(0u, pProgram->getNumKernels());
}
HWTEST_P(ProgramFromBinaryTest, givenProgramWhenCleanCurrentKernelInfoIsCalledButGpuIsNotYetDoneThenKernelAllocationIsPutOnDeferredFreeListAndCsrRegistersCacheFlush) {
cl_device_id device = pClDevice;
auto &csr = pDevice->getGpgpuCommandStreamReceiver();
EXPECT_TRUE(csr.getTemporaryAllocations().peekIsEmpty());
pProgram->build(1, &device, nullptr, true);
pProgram->build(pProgram->getDevices(), nullptr, true);
auto kernelAllocation = pProgram->getKernelInfo(size_t(0))->getGraphicsAllocation();
kernelAllocation->updateTaskCount(100, csr.getOsContext().getContextId());
*csr.getTagAddress() = 0;
@ -652,9 +642,7 @@ HWTEST_P(ProgramFromBinaryTest, givenIsaAllocationUsedByMultipleCsrsWhenItIsDele
auto &csr0 = this->pDevice->getUltCommandStreamReceiverFromIndex<FamilyType>(0u);
auto &csr1 = this->pDevice->getUltCommandStreamReceiverFromIndex<FamilyType>(1u);
cl_device_id device = pClDevice;
pProgram->build(1, &device, nullptr, true);
pProgram->build(pProgram->getDevices(), nullptr, true);
auto kernelAllocation = pProgram->getKernelInfo(size_t(0))->getGraphicsAllocation();
@ -679,13 +667,8 @@ TEST_P(ProgramFromSourceTest, GivenSpecificParamatersWhenBuildingProgramThenSucc
KernelBinaryHelper kbHelper(BinaryFileName, true);
auto device = pPlatform->getClDevice(0);
cl_device_id deviceList = {0};
cl_device_id usedDevice = pPlatform->getClDevice(0);
CreateProgramWithSource(
pContext,
&usedDevice,
SourceFileName);
// Order of following microtests is important - do not change.
@ -693,20 +676,9 @@ TEST_P(ProgramFromSourceTest, GivenSpecificParamatersWhenBuildingProgramThenSucc
auto pMockProgram = pProgram;
// invalid build parameters: combinations of numDevices & deviceList
retVal = pProgram->build(1, nullptr, nullptr, false);
EXPECT_EQ(CL_INVALID_VALUE, retVal);
retVal = pProgram->build(0, &deviceList, nullptr, false);
EXPECT_EQ(CL_INVALID_VALUE, retVal);
// invalid build parameters: invalid content of deviceList
retVal = pProgram->build(1, &deviceList, nullptr, false);
EXPECT_EQ(CL_INVALID_DEVICE, retVal);
// fail build - another build is already in progress
pMockProgram->setBuildStatus(CL_BUILD_IN_PROGRESS);
retVal = pProgram->build(0, nullptr, nullptr, false);
retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
EXPECT_EQ(CL_INVALID_OPERATION, retVal);
pMockProgram->setBuildStatus(CL_BUILD_NONE);
@ -716,13 +688,13 @@ TEST_P(ProgramFromSourceTest, GivenSpecificParamatersWhenBuildingProgramThenSucc
std::unique_ptr<RootDeviceEnvironment> rootDeviceEnvironment = std::make_unique<NoCompilerInterfaceRootDeviceEnvironment>(*executionEnvironment);
std::swap(rootDeviceEnvironment, executionEnvironment->rootDeviceEnvironments[device->getRootDeviceIndex()]);
auto p2 = std::make_unique<MockProgram>(toClDeviceVector(*device));
retVal = p2->build(0, nullptr, nullptr, false);
retVal = p2->build(p2->getDevices(), nullptr, false);
EXPECT_EQ(CL_OUT_OF_HOST_MEMORY, retVal);
p2.reset(nullptr);
std::swap(rootDeviceEnvironment, executionEnvironment->rootDeviceEnvironments[device->getRootDeviceIndex()]);
// fail build - any build error (here caused by specifying unrecognized option)
retVal = pProgram->build(0, nullptr, "-invalid-option", false);
retVal = pProgram->build(pProgram->getDevices(), "-invalid-option", false);
EXPECT_EQ(CL_BUILD_PROGRAM_FAILURE, retVal);
// fail build - linked code is corrupted and cannot be postprocessed
@ -736,15 +708,21 @@ TEST_P(ProgramFromSourceTest, GivenSpecificParamatersWhenBuildingProgramThenSucc
EXPECT_NE(nullptr, pSourceBuffer);
p3->sourceCode = pSourceBuffer.get();
p3->createdFrom = Program::CreatedFrom::SOURCE;
retVal = p3->build(0, nullptr, nullptr, false);
retVal = p3->build(p3->getDevices(), nullptr, false);
EXPECT_EQ(CL_INVALID_BINARY, retVal);
p3.reset(nullptr);
// build successfully - build kernel and write it to Kernel Cache
pMockProgram->clearOptions();
retVal = pProgram->build(0, nullptr, nullptr, false);
std::string receivedInternalOptions;
auto debugVars = NEO::getFclDebugVars();
debugVars.receivedInternalOptionsOutput = &receivedInternalOptions;
gEnvironment->fclPushDebugVars(debugVars);
retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_TRUE(CompilerOptions::contains(pProgram->getInternalOptions(), pPlatform->getClDevice(0)->peekCompilerExtensions())) << pProgram->getInternalOptions();
EXPECT_TRUE(CompilerOptions::contains(receivedInternalOptions, pPlatform->getClDevice(0)->peekCompilerExtensions())) << receivedInternalOptions;
gEnvironment->fclPopDebugVars();
// get build log
size_t param_value_size_ret = 0u;
@ -770,11 +748,11 @@ TEST_P(ProgramFromSourceTest, GivenSpecificParamatersWhenBuildingProgramThenSucc
// build successfully - build kernel but do not write it to Kernel Cache (kernel is already in the Cache)
pMockProgram->setBuildStatus(CL_BUILD_NONE);
retVal = pProgram->build(0, nullptr, nullptr, false);
retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
EXPECT_EQ(CL_SUCCESS, retVal);
// build successfully - kernel is already in Kernel Cache, do not build and take it from Cache
retVal = pProgram->build(0, nullptr, nullptr, true);
retVal = pProgram->build(pProgram->getDevices(), nullptr, true);
EXPECT_EQ(CL_SUCCESS, retVal);
// fail build - code to be build does not exist
@ -782,41 +760,41 @@ TEST_P(ProgramFromSourceTest, GivenSpecificParamatersWhenBuildingProgramThenSucc
pMockProgram->createdFrom = Program::CreatedFrom::SOURCE;
pMockProgram->setBuildStatus(CL_BUILD_NONE);
pMockProgram->setCreatedFromBinary(false);
retVal = pProgram->build(0, nullptr, nullptr, false);
retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
EXPECT_EQ(CL_INVALID_PROGRAM, retVal);
}
TEST_P(ProgramFromSourceTest, CreateWithSource_Build_Options_Duplicate) {
KernelBinaryHelper kbHelper(BinaryFileName, false);
retVal = pProgram->build(0, nullptr, nullptr, false);
retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
EXPECT_EQ(CL_SUCCESS, retVal);
retVal = pProgram->build(0, nullptr, CompilerOptions::fastRelaxedMath.data(), false);
retVal = pProgram->build(pProgram->getDevices(), CompilerOptions::fastRelaxedMath.data(), false);
EXPECT_EQ(CL_SUCCESS, retVal);
retVal = pProgram->build(0, nullptr, CompilerOptions::fastRelaxedMath.data(), false);
retVal = pProgram->build(pProgram->getDevices(), CompilerOptions::fastRelaxedMath.data(), false);
EXPECT_EQ(CL_SUCCESS, retVal);
retVal = pProgram->build(0, nullptr, CompilerOptions::finiteMathOnly.data(), false);
retVal = pProgram->build(pProgram->getDevices(), CompilerOptions::finiteMathOnly.data(), false);
EXPECT_EQ(CL_SUCCESS, retVal);
retVal = pProgram->build(0, nullptr, nullptr, false);
retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
EXPECT_EQ(CL_SUCCESS, retVal);
}
TEST_P(ProgramFromSourceTest, WhenBuildingProgramThenFeaturesOptionIsNotAdded) {
auto cip = new MockCompilerInterfaceCaptureBuildOptions();
auto pClDevice = pContext->getDevice(0);
pClDevice->getExecutionEnvironment()->rootDeviceEnvironments[pClDevice->getRootDeviceIndex()]->compilerInterface.reset(cip);
auto featuresOption = static_cast<ClDevice *>(devices[0])->peekCompilerFeatures();
EXPECT_THAT(pProgram->getInternalOptions(), testing::Not(testing::HasSubstr(featuresOption)));
retVal = pProgram->build(1, devices, nullptr, false);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_THAT(pProgram->getInternalOptions(), testing::Not(testing::HasSubstr(featuresOption)));
retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
EXPECT_THAT(cip->buildInternalOptions, testing::Not(testing::HasSubstr(featuresOption)));
}
TEST_P(ProgramFromSourceTest, WhenBuildingProgramWithOpenClC30ThenFeaturesOptionIsAdded) {
auto featuresOption = static_cast<ClDevice *>(devices[0])->peekCompilerFeatures();
EXPECT_THAT(pProgram->getInternalOptions(), testing::Not(testing::HasSubstr(featuresOption)));
auto cip = new MockCompilerInterfaceCaptureBuildOptions();
auto pClDevice = pContext->getDevice(0);
@ -825,9 +803,12 @@ TEST_P(ProgramFromSourceTest, WhenBuildingProgramWithOpenClC30ThenFeaturesOption
pProgram->sourceCode = "__kernel mock() {}";
pProgram->createdFrom = Program::CreatedFrom::SOURCE;
retVal = pProgram->build(1, devices, "-cl-std=CL3.0", false);
MockProgram::initInternalOptionsCalled = 0;
retVal = pProgram->build(pProgram->getDevices(), "-cl-std=CL3.0", false);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_THAT(pProgram->getInternalOptions(), testing::HasSubstr(featuresOption));
EXPECT_THAT(cip->buildInternalOptions, testing::HasSubstr(featuresOption));
EXPECT_EQ(1, MockProgram::initInternalOptionsCalled);
}
TEST_P(ProgramFromSourceTest, WhenBuildingProgramWithOpenClC30ThenFeaturesOptionIsAddedOnlyOnce) {
@ -838,13 +819,13 @@ TEST_P(ProgramFromSourceTest, WhenBuildingProgramWithOpenClC30ThenFeaturesOption
pProgram->sourceCode = "__kernel mock() {}";
pProgram->createdFrom = Program::CreatedFrom::SOURCE;
retVal = pProgram->build(0, nullptr, "-cl-std=CL3.0", false);
retVal = pProgram->build(pProgram->getDevices(), "-cl-std=CL3.0", false);
EXPECT_EQ(CL_SUCCESS, retVal);
retVal = pProgram->build(0, nullptr, "-cl-std=CL3.0", false);
retVal = pProgram->build(pProgram->getDevices(), "-cl-std=CL3.0", false);
EXPECT_EQ(CL_SUCCESS, retVal);
auto expectedFeaturesOption = static_cast<ClDevice *>(devices[0])->peekCompilerFeatures();
auto &internalOptions = pProgram->getInternalOptions();
auto &internalOptions = cip->buildInternalOptions;
auto pos = internalOptions.find(expectedFeaturesOption);
EXPECT_NE(std::string::npos, pos);
@ -859,9 +840,11 @@ TEST_P(ProgramFromSourceTest, WhenCompilingProgramThenFeaturesOptionIsNotAdded)
auto featuresOption = pClDevice->peekCompilerFeatures();
EXPECT_THAT(pCompilerInterface->buildInternalOptions, testing::Not(testing::HasSubstr(featuresOption)));
MockProgram::initInternalOptionsCalled = 0;
retVal = pProgram->compile(pProgram->getDevices(), nullptr, 0, nullptr, nullptr);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_THAT(pCompilerInterface->buildInternalOptions, testing::Not(testing::HasSubstr(featuresOption)));
EXPECT_EQ(1, MockProgram::initInternalOptionsCalled);
}
TEST_P(ProgramFromSourceTest, WhenCompilingProgramWithOpenClC30ThenFeaturesOptionIsAdded) {
@ -911,22 +894,20 @@ std::map<const void *, uint32_t> Callback::watchList;
TEST_P(ProgramFromSourceTest, GivenDifferentCommpilerOptionsWhenBuildingProgramThenKernelHashesAreDifferent) {
KernelBinaryHelper kbHelper(BinaryFileName, true);
cl_device_id usedDevice = pPlatform->getClDevice(0);
CreateProgramWithSource(
pContext,
&usedDevice,
SourceFileName);
Callback callback;
retVal = pProgram->build(0, nullptr, nullptr, true);
retVal = pProgram->build(pProgram->getDevices(), nullptr, true);
EXPECT_EQ(CL_SUCCESS, retVal);
auto hash1 = pProgram->getCachedFileName();
auto kernel1 = pProgram->getKernelInfo("CopyBuffer");
Callback::watch(kernel1);
EXPECT_NE(nullptr, kernel1);
retVal = pProgram->build(0, nullptr, CompilerOptions::fastRelaxedMath.data(), true);
retVal = pProgram->build(pProgram->getDevices(), CompilerOptions::fastRelaxedMath.data(), true);
EXPECT_EQ(CL_SUCCESS, retVal);
auto hash2 = pProgram->getCachedFileName();
auto kernel2 = pProgram->getKernelInfo("CopyBuffer");
@ -935,7 +916,7 @@ TEST_P(ProgramFromSourceTest, GivenDifferentCommpilerOptionsWhenBuildingProgramT
Callback::unwatch(kernel1);
Callback::watch(kernel2);
retVal = pProgram->build(0, nullptr, CompilerOptions::finiteMathOnly.data(), true);
retVal = pProgram->build(pProgram->getDevices(), CompilerOptions::finiteMathOnly.data(), true);
EXPECT_EQ(CL_SUCCESS, retVal);
auto hash3 = pProgram->getCachedFileName();
auto kernel3 = pProgram->getKernelInfo("CopyBuffer");
@ -948,7 +929,7 @@ TEST_P(ProgramFromSourceTest, GivenDifferentCommpilerOptionsWhenBuildingProgramT
pProgram->createdFrom = NEO::Program::CreatedFrom::BINARY;
pProgram->setIrBinary(new char[16], true);
pProgram->setIrBinarySize(16, true);
retVal = pProgram->build(0, nullptr, nullptr, true);
retVal = pProgram->build(pProgram->getDevices(), nullptr, true);
EXPECT_EQ(CL_SUCCESS, retVal);
auto hash4 = pProgram->getCachedFileName();
auto kernel4 = pProgram->getKernelInfo("CopyBuffer");
@ -958,7 +939,7 @@ TEST_P(ProgramFromSourceTest, GivenDifferentCommpilerOptionsWhenBuildingProgramT
Callback::watch(kernel4);
pProgram->createdFrom = NEO::Program::CreatedFrom::SOURCE;
retVal = pProgram->build(0, nullptr, nullptr, true);
retVal = pProgram->build(pProgram->getDevices(), nullptr, true);
EXPECT_EQ(CL_SUCCESS, retVal);
auto hash5 = pProgram->getCachedFileName();
auto kernel5 = pProgram->getKernelInfo("CopyBuffer");
@ -975,11 +956,8 @@ TEST_P(ProgramFromSourceTest, GivenEmptyProgramWhenCreatingProgramThenInvalidVal
}
TEST_P(ProgramFromSourceTest, GivenSpecificParamatersWhenCompilingProgramThenSuccessOrCorrectErrorCodeIsReturned) {
cl_device_id usedDevice = pPlatform->getClDevice(0);
CreateProgramWithSource(
pContext,
&usedDevice,
SourceFileName);
cl_program inputHeaders;
@ -1101,12 +1079,14 @@ TEST_F(ProgramTests, GivenFlagsWhenLinkingProgramThenBuildOptionsHaveBeenApplied
auto pProgram = std::make_unique<SucceedingGenBinaryProgram>(toClDeviceVector(*pClDevice));
pProgram->sourceCode = "__kernel mock() {}";
pProgram->createdFrom = Program::CreatedFrom::SOURCE;
MockProgram::initInternalOptionsCalled = 0;
cl_program program = pProgram.get();
// compile successfully a kernel to be linked later
cl_int retVal = pProgram->compile(pProgram->getDevices(), nullptr, 0, nullptr, nullptr);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_EQ(1, MockProgram::initInternalOptionsCalled);
// Ask to link created program with NEO::CompilerOptions::gtpinRera and NEO::CompilerOptions::greaterThan4gbBuffersRequired flags.
auto options = CompilerOptions::concatenate(CompilerOptions::greaterThan4gbBuffersRequired, CompilerOptions::gtpinRera, CompilerOptions::finiteMathOnly);
@ -1115,6 +1095,7 @@ TEST_F(ProgramTests, GivenFlagsWhenLinkingProgramThenBuildOptionsHaveBeenApplied
retVal = pProgram->link(pProgram->getDevices(), options.c_str(), 1, &program);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_EQ(2, MockProgram::initInternalOptionsCalled);
// Check build options that were applied
EXPECT_FALSE(CompilerOptions::contains(cip->buildOptions, CompilerOptions::fastRelaxedMath)) << cip->buildOptions;
@ -1176,10 +1157,8 @@ TEST_P(ProgramFromSourceTest, GivenAdvancedOptionsWhenCreatingProgramThenSuccess
}
TEST_P(ProgramFromSourceTest, GivenSpecificParamatersWhenLinkingProgramThenSuccessOrCorrectErrorCodeIsReturned) {
cl_device_id usedDevice = pPlatform->getClDevice(0);
CreateProgramWithSource(
pContext,
&usedDevice,
SourceFileName);
cl_program program = pProgram;
@ -1235,8 +1214,7 @@ TEST_P(ProgramFromSourceTest, GivenSpecificParamatersWhenLinkingProgramThenSucce
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));
auto p2 = std::make_unique<FailingGenBinaryProgram>(pProgram->getDevices());
retVal = p2->link(p2->getDevices(), nullptr, 1, &program);
EXPECT_EQ(CL_INVALID_BINARY, retVal);
p2.reset(nullptr);
@ -1307,14 +1285,11 @@ class CommandStreamReceiverMock : public UltCommandStreamReceiver<FamilyType> {
};
HWTEST_F(PatchTokenTests, givenKernelRequiringConstantAllocationWhenMakeResidentIsCalledThenConstantAllocationIsMadeResident) {
cl_device_id device = pClDevice;
CreateProgramFromBinary(pContext, pContext->getDevices(), "test_constant_memory");
ASSERT_NE(nullptr, pProgram);
retVal = pProgram->build(
1,
&device,
pProgram->getDevices(),
nullptr,
false);
@ -1374,14 +1349,11 @@ HWTEST_F(PatchTokenTests, givenKernelRequiringConstantAllocationWhenMakeResident
}
TEST_F(PatchTokenTests, WhenBuildingProgramThenGwsIsSet) {
cl_device_id device = pClDevice;
CreateProgramFromBinary(pContext, pContext->getDevices(), "kernel_data_param");
ASSERT_NE(nullptr, pProgram);
retVal = pProgram->build(
1,
&device,
pProgram->getDevices(),
nullptr,
false);
@ -1396,14 +1368,11 @@ TEST_F(PatchTokenTests, WhenBuildingProgramThenGwsIsSet) {
}
TEST_F(PatchTokenTests, WhenBuildingProgramThenLwsIsSet) {
cl_device_id device = pClDevice;
CreateProgramFromBinary(pContext, pContext->getDevices(), "kernel_data_param");
ASSERT_NE(nullptr, pProgram);
retVal = pProgram->build(
1,
&device,
pProgram->getDevices(),
nullptr,
false);
@ -1429,14 +1398,12 @@ TEST_F(PatchTokenTests, WhenBuildingProgramThenLwsIsSet) {
TEST_F(PatchTokenTests, WhenBuildingProgramThenConstantKernelArgsAreAvailable) {
// PATCH_TOKEN_STATELESS_CONSTANT_MEMORY_OBJECT_KERNEL_ARGUMENT
cl_device_id device = pClDevice;
CreateProgramFromBinary(pContext, pContext->getDevices(), "test_basic_constant");
ASSERT_NE(nullptr, pProgram);
retVal = pProgram->build(
1,
&device,
pProgram->getDevices(),
nullptr,
false);
@ -1470,14 +1437,12 @@ TEST_F(PatchTokenTests, GivenVmeKernelWhenBuildingKernelThenArgAvailable) {
GTEST_SKIP();
}
// PATCH_TOKEN_INLINE_VME_SAMPLER_INFO token indicates a VME kernel.
cl_device_id device = pClDevice;
CreateProgramFromBinary(pContext, pContext->getDevices(), "vme_kernels");
ASSERT_NE(nullptr, pProgram);
retVal = pProgram->build(
1,
&device,
pProgram->getDevices(),
nullptr,
false);
@ -1584,15 +1549,12 @@ INSTANTIATE_TEST_CASE_P(ProgramFromSourceTests,
using ProgramWithDebugSymbolsTests = Test<ProgramSimpleFixture>;
TEST_F(ProgramWithDebugSymbolsTests, GivenProgramCreatedWithDashGOptionWhenGettingProgramBinariesThenDebugDataIsIncluded) {
cl_device_id device = pClDevice;
CreateProgramFromBinary(pContext, pContext->getDevices(), "CopyBuffer_simd16", "-g");
ASSERT_NE(nullptr, pProgram);
retVal = pProgram->build(
1,
&device,
pProgram->getDevices(),
"-g",
false);
EXPECT_EQ(CL_SUCCESS, retVal);
@ -1646,12 +1608,13 @@ TEST_F(ProgramTests, WhenProgramIsCreatedThenCorrectOclVersionIsInOptions) {
DebugManager.flags.DisableStatelessToStatefulOptimization.set(false);
MockProgram program(pContext, false, toClDeviceVector(*pClDevice));
auto internalOptions = program.getInitInternalOptions();
if (pClDevice->getEnabledClVersion() == 30) {
EXPECT_TRUE(CompilerOptions::contains(program.getInternalOptions(), "-ocl-version=300")) << program.getInternalOptions();
EXPECT_TRUE(CompilerOptions::contains(internalOptions, "-ocl-version=300")) << internalOptions;
} else if (pClDevice->getEnabledClVersion() == 21) {
EXPECT_TRUE(CompilerOptions::contains(program.getInternalOptions(), "-ocl-version=210")) << program.getInternalOptions();
EXPECT_TRUE(CompilerOptions::contains(internalOptions, "-ocl-version=210")) << internalOptions;
} else {
EXPECT_TRUE(CompilerOptions::contains(program.getInternalOptions(), "-ocl-version=120")) << program.getInternalOptions();
EXPECT_TRUE(CompilerOptions::contains(internalOptions, "-ocl-version=120")) << internalOptions;
}
}
@ -1665,7 +1628,8 @@ TEST_F(ProgramTests, GivenForcedClVersionWhenProgramIsCreatedThenCorrectOclOptio
for (auto &testedValue : testedValues) {
pClDevice->enabledClVersion = testedValue.first;
MockProgram program{pContext, false, toClDeviceVector(*pClDevice)};
EXPECT_TRUE(CompilerOptions::contains(program.getInternalOptions(), testedValue.second));
auto internalOptions = program.getInitInternalOptions();
EXPECT_TRUE(CompilerOptions::contains(internalOptions, testedValue.second));
}
}
@ -1674,7 +1638,8 @@ TEST_F(ProgramTests, GivenStatelessToStatefulIsDisabledWhenProgramIsCreatedThenG
DebugManager.flags.DisableStatelessToStatefulOptimization.set(true);
MockProgram program(pContext, false, toClDeviceVector(*pClDevice));
EXPECT_TRUE(CompilerOptions::contains(program.getInternalOptions(), NEO::CompilerOptions::greaterThan4gbBuffersRequired));
auto internalOptions = program.getInitInternalOptions();
EXPECT_TRUE(CompilerOptions::contains(internalOptions, NEO::CompilerOptions::greaterThan4gbBuffersRequired));
}
TEST_F(ProgramTests, WhenCreatingProgramThenBindlessIsEnabledOnlyIfDebugFlagIsEnabled) {
@ -1685,15 +1650,17 @@ TEST_F(ProgramTests, WhenCreatingProgramThenBindlessIsEnabledOnlyIfDebugFlagIsEn
DebugManager.flags.UseBindlessMode.set(0);
MockProgram programNoBindless(pContext, false, toClDeviceVector(*pClDevice));
EXPECT_FALSE(CompilerOptions::contains(programNoBindless.getInternalOptions(), CompilerOptions::bindlessBuffers)) << programNoBindless.getInternalOptions();
EXPECT_FALSE(CompilerOptions::contains(programNoBindless.getInternalOptions(), CompilerOptions::bindlessImages)) << programNoBindless.getInternalOptions();
auto internalOptionsNoBindless = programNoBindless.getInitInternalOptions();
EXPECT_FALSE(CompilerOptions::contains(internalOptionsNoBindless, CompilerOptions::bindlessBuffers)) << internalOptionsNoBindless;
EXPECT_FALSE(CompilerOptions::contains(internalOptionsNoBindless, CompilerOptions::bindlessImages)) << internalOptionsNoBindless;
}
{
DebugManager.flags.UseBindlessMode.set(1);
MockProgram programNoBindless(pContext, false, toClDeviceVector(*pClDevice));
EXPECT_TRUE(CompilerOptions::contains(programNoBindless.getInternalOptions(), CompilerOptions::bindlessBuffers)) << programNoBindless.getInternalOptions();
EXPECT_TRUE(CompilerOptions::contains(programNoBindless.getInternalOptions(), CompilerOptions::bindlessImages)) << programNoBindless.getInternalOptions();
MockProgram programBindless(pContext, false, toClDeviceVector(*pClDevice));
auto internalOptionsBindless = programBindless.getInitInternalOptions();
EXPECT_TRUE(CompilerOptions::contains(internalOptionsBindless, CompilerOptions::bindlessBuffers)) << internalOptionsBindless;
EXPECT_TRUE(CompilerOptions::contains(internalOptionsBindless, CompilerOptions::bindlessImages)) << internalOptionsBindless;
}
}
@ -1701,7 +1668,8 @@ TEST_F(ProgramTests, givenDeviceThatSupportsSharedSystemMemoryAllocationWhenProg
pClDevice->deviceInfo.sharedSystemMemCapabilities = CL_UNIFIED_SHARED_MEMORY_ACCESS_INTEL | CL_UNIFIED_SHARED_MEMORY_ATOMIC_ACCESS_INTEL | CL_UNIFIED_SHARED_MEMORY_CONCURRENT_ACCESS_INTEL | CL_UNIFIED_SHARED_MEMORY_CONCURRENT_ATOMIC_ACCESS_INTEL;
pClDevice->sharedDeviceInfo.sharedSystemAllocationsSupport = true;
MockProgram program(pContext, false, toClDeviceVector(*pClDevice));
EXPECT_TRUE(CompilerOptions::contains(program.getInternalOptions().c_str(), CompilerOptions::greaterThan4gbBuffersRequired)) << program.getInternalOptions();
auto internalOptions = program.getInitInternalOptions();
EXPECT_TRUE(CompilerOptions::contains(internalOptions.c_str(), CompilerOptions::greaterThan4gbBuffersRequired)) << internalOptions;
}
TEST_F(ProgramTests, GivenForce32BitAddressessWhenProgramIsCreatedThenGreaterThan4gbBuffersRequiredIsCorrectlySet) {
@ -1712,10 +1680,11 @@ TEST_F(ProgramTests, GivenForce32BitAddressessWhenProgramIsCreatedThenGreaterTha
if (pDevice) {
const_cast<DeviceInfo *>(&pDevice->getDeviceInfo())->force32BitAddressess = true;
MockProgram program(pContext, false, toClDeviceVector(*pClDevice));
auto internalOptions = program.getInitInternalOptions();
if (pDevice->areSharedSystemAllocationsAllowed()) {
EXPECT_TRUE(CompilerOptions::contains(program.getInternalOptions(), CompilerOptions::greaterThan4gbBuffersRequired)) << program.getInternalOptions();
EXPECT_TRUE(CompilerOptions::contains(internalOptions, CompilerOptions::greaterThan4gbBuffersRequired)) << internalOptions;
} else {
EXPECT_FALSE(CompilerOptions::contains(program.getInternalOptions(), NEO::CompilerOptions::greaterThan4gbBuffersRequired)) << program.getInternalOptions();
EXPECT_FALSE(CompilerOptions::contains(internalOptions, NEO::CompilerOptions::greaterThan4gbBuffersRequired)) << internalOptions;
}
} else {
EXPECT_NE(CL_DEVICE_NOT_FOUND, retVal);
@ -1728,10 +1697,11 @@ TEST_F(ProgramTests, Given32bitSupportWhenProgramIsCreatedThenGreaterThan4gbBuff
DebugManager.flags.DisableStatelessToStatefulOptimization.set(false);
std::unique_ptr<MockProgram> program{Program::createBuiltInFromSource<MockProgram>("", pContext, pContext->getDevices(), nullptr)};
auto internalOptions = program->getInitInternalOptions();
if ((false == pDevice->areSharedSystemAllocationsAllowed()) && (false == is32bit)) {
EXPECT_FALSE(CompilerOptions::contains(program->getInternalOptions(), NEO::CompilerOptions::greaterThan4gbBuffersRequired)) << program->getInternalOptions();
EXPECT_FALSE(CompilerOptions::contains(internalOptions, NEO::CompilerOptions::greaterThan4gbBuffersRequired)) << internalOptions;
} else {
EXPECT_TRUE(CompilerOptions::contains(program->getInternalOptions(), NEO::CompilerOptions::greaterThan4gbBuffersRequired)) << program->getInternalOptions();
EXPECT_TRUE(CompilerOptions::contains(internalOptions, NEO::CompilerOptions::greaterThan4gbBuffersRequired)) << internalOptions;
}
DebugManager.flags.DisableStatelessToStatefulOptimization.set(defaultSetting);
}
@ -1741,13 +1711,15 @@ TEST_F(ProgramTests, GivenStatelessToStatefulIsDisabledWhenProgramIsCreatedThenG
DebugManager.flags.DisableStatelessToStatefulOptimization.set(true);
std::unique_ptr<MockProgram> program{Program::createBuiltInFromSource<MockProgram>("", pContext, pContext->getDevices(), nullptr)};
EXPECT_TRUE(CompilerOptions::contains(program->getInternalOptions(), NEO::CompilerOptions::greaterThan4gbBuffersRequired)) << program->getInternalOptions();
auto internalOptions = program->getInitInternalOptions();
EXPECT_TRUE(CompilerOptions::contains(internalOptions, NEO::CompilerOptions::greaterThan4gbBuffersRequired)) << internalOptions;
DebugManager.flags.DisableStatelessToStatefulOptimization.set(defaultSetting);
}
TEST_F(ProgramTests, givenProgramWhenItIsCompiledThenItAlwaysHavePreserveVec3TypeInternalOptionSet) {
std::unique_ptr<MockProgram> program(Program::createBuiltInFromSource<MockProgram>("", pContext, pContext->getDevices(), nullptr));
EXPECT_TRUE(CompilerOptions::contains(program->getInternalOptions(), CompilerOptions::preserveVec3Type)) << program->getInternalOptions();
auto internalOptions = program->getInitInternalOptions();
EXPECT_TRUE(CompilerOptions::contains(internalOptions, CompilerOptions::preserveVec3Type)) << internalOptions;
}
TEST_F(ProgramTests, Force32BitAddressessWhenProgramIsCreatedThenGreaterThan4gbBuffersRequiredIsCorrectlySet) {
@ -1756,13 +1728,14 @@ TEST_F(ProgramTests, Force32BitAddressessWhenProgramIsCreatedThenGreaterThan4gbB
DebugManager.flags.DisableStatelessToStatefulOptimization.set(false);
const_cast<DeviceInfo *>(&pDevice->getDeviceInfo())->force32BitAddressess = true;
std::unique_ptr<MockProgram> program{Program::createBuiltInFromSource<MockProgram>("", pContext, pContext->getDevices(), nullptr)};
auto internalOptions = program->getInitInternalOptions();
if (is32bit) {
EXPECT_TRUE(CompilerOptions::contains(program->getInternalOptions(), CompilerOptions::greaterThan4gbBuffersRequired)) << program->getInternalOptions();
EXPECT_TRUE(CompilerOptions::contains(internalOptions, CompilerOptions::greaterThan4gbBuffersRequired)) << internalOptions;
} else {
if (false == pDevice->areSharedSystemAllocationsAllowed()) {
EXPECT_FALSE(CompilerOptions::contains(program->getInternalOptions(), NEO::CompilerOptions::greaterThan4gbBuffersRequired)) << program->getInternalOptions();
EXPECT_FALSE(CompilerOptions::contains(internalOptions, NEO::CompilerOptions::greaterThan4gbBuffersRequired)) << internalOptions;
} else {
EXPECT_TRUE(CompilerOptions::contains(program->getInternalOptions(), NEO::CompilerOptions::greaterThan4gbBuffersRequired)) << program->getInternalOptions();
EXPECT_TRUE(CompilerOptions::contains(internalOptions, NEO::CompilerOptions::greaterThan4gbBuffersRequired)) << internalOptions;
}
}
DebugManager.flags.DisableStatelessToStatefulOptimization.set(defaultSetting);
@ -1777,8 +1750,9 @@ TEST_F(ProgramTests, GivenStatelessToStatefulBufferOffsetOptimizationWhenProgram
const char **programSources = reinterpret_cast<const char **>(&programPointer);
size_t length = sizeof(programSource);
std::unique_ptr<MockProgram> program(Program::create<MockProgram>(pContext, 1u, programSources, &length, errorCode));
auto internalOptions = program->getInitInternalOptions();
EXPECT_TRUE(CompilerOptions::contains(program->getInternalOptions(), CompilerOptions::hasBufferOffsetArg)) << program->getInternalOptions();
EXPECT_TRUE(CompilerOptions::contains(internalOptions, CompilerOptions::hasBufferOffsetArg)) << internalOptions;
}
TEST_F(ProgramTests, givenStatelessToStatefullOptimizationOffWHenProgramIsCreatedThenOptimizationStringIsNotPresent) {
@ -1790,7 +1764,8 @@ TEST_F(ProgramTests, givenStatelessToStatefullOptimizationOffWHenProgramIsCreate
const char **programSources = reinterpret_cast<const char **>(&programPointer);
size_t length = sizeof(programSource);
std::unique_ptr<MockProgram> program(Program::create<MockProgram>(pContext, 1u, programSources, &length, errorCode));
EXPECT_FALSE(CompilerOptions::contains(program->getInternalOptions(), CompilerOptions::hasBufferOffsetArg)) << program->getInternalOptions();
auto internalOptions = program->getInitInternalOptions();
EXPECT_FALSE(CompilerOptions::contains(internalOptions, CompilerOptions::hasBufferOffsetArg)) << internalOptions;
}
TEST_F(ProgramTests, GivenContextWhenCreateProgramThenIncrementContextRefCount) {
@ -1966,7 +1941,7 @@ TEST_F(ProgramTests, GivenGtpinReraFlagWhenBuildingProgramThenCorrectOptionsAreS
program->createdFrom = Program::CreatedFrom::SOURCE;
// Ask to build created program without NEO::CompilerOptions::gtpinRera flag.
cl_int retVal = program->build(0, nullptr, CompilerOptions::fastRelaxedMath.data(), false);
cl_int retVal = program->build(program->getDevices(), CompilerOptions::fastRelaxedMath.data(), false);
EXPECT_EQ(CL_SUCCESS, retVal);
// Check build options that were applied
@ -1976,7 +1951,7 @@ TEST_F(ProgramTests, GivenGtpinReraFlagWhenBuildingProgramThenCorrectOptionsAreS
// Ask to build created program with NEO::CompilerOptions::gtpinRera flag.
cip->buildOptions.clear();
cip->buildInternalOptions.clear();
retVal = program->build(0, nullptr, CompilerOptions::concatenate(CompilerOptions::gtpinRera, CompilerOptions::finiteMathOnly).c_str(), false);
retVal = program->build(program->getDevices(), CompilerOptions::concatenate(CompilerOptions::gtpinRera, CompilerOptions::finiteMathOnly).c_str(), false);
EXPECT_EQ(CL_SUCCESS, retVal);
// Check build options that were applied
@ -2137,13 +2112,13 @@ class Program32BitTests : public ProgramTests {
TEST_F(Program32BitTests, givenDeviceWithForce32BitAddressingOnWhenBuiltinIsCreatedThenNoFlagsArePassedAsInternalOptions) {
MockProgram program(toClDeviceVector(*pClDevice));
auto &internalOptions = program.getInternalOptions();
auto internalOptions = program.getInitInternalOptions();
EXPECT_THAT(internalOptions, testing::HasSubstr(std::string("")));
}
TEST_F(Program32BitTests, givenDeviceWithForce32BitAddressingOnWhenProgramIsCreatedThen32bitFlagIsPassedAsInternalOption) {
MockProgram program(pContext, false, toClDeviceVector(*pClDevice));
auto &internalOptions = program.getInternalOptions();
auto internalOptions = program.getInitInternalOptions();
std::string s1 = internalOptions;
size_t pos = s1.find(NEO::CompilerOptions::arch32bit.data());
if (is64bit) {
@ -2154,15 +2129,15 @@ TEST_F(Program32BitTests, givenDeviceWithForce32BitAddressingOnWhenProgramIsCrea
}
TEST_F(ProgramTests, givenNewProgramTheStatelessToStatefulBufferOffsetOtimizationIsMatchingThePlatformEnablingStatus) {
MockProgram prog(pContext, false, toClDeviceVector(*pClDevice));
auto &internalOpts = prog.getInternalOptions();
MockProgram program(pContext, false, toClDeviceVector(*pClDevice));
auto internalOptions = program.getInitInternalOptions();
HardwareCapabilities hwCaps = {0};
HwHelper::get(prog.getDevice().getHardwareInfo().platform.eRenderCoreFamily).setupHardwareCapabilities(&hwCaps, prog.getDevice().getHardwareInfo());
HwHelper::get(program.getDevice().getHardwareInfo().platform.eRenderCoreFamily).setupHardwareCapabilities(&hwCaps, program.getDevice().getHardwareInfo());
if (hwCaps.isStatelesToStatefullWithOffsetSupported) {
EXPECT_TRUE(CompilerOptions::contains(internalOpts, CompilerOptions::hasBufferOffsetArg));
EXPECT_TRUE(CompilerOptions::contains(internalOptions, CompilerOptions::hasBufferOffsetArg));
} else {
EXPECT_FALSE(CompilerOptions::contains(internalOpts, CompilerOptions::hasBufferOffsetArg));
EXPECT_FALSE(CompilerOptions::contains(internalOptions, CompilerOptions::hasBufferOffsetArg));
}
}
@ -2536,39 +2511,40 @@ TEST_F(ProgramTests, whenRebuildingProgramThenStoreDeviceBinaryProperly) {
TEST_F(ProgramTests, givenProgramWhenInternalOptionsArePassedThenTheyAreAddedToProgramInternalOptions) {
MockProgram program(toClDeviceVector(*pClDevice));
program.getInternalOptions().erase();
std::string buildOptions = NEO::CompilerOptions::gtpinRera.str();
program.extractInternalOptions(buildOptions);
EXPECT_STREQ(program.getInternalOptions().c_str(), NEO::CompilerOptions::gtpinRera.data());
std::string internalOptions;
program.extractInternalOptions(buildOptions, internalOptions);
EXPECT_STREQ(internalOptions.c_str(), NEO::CompilerOptions::gtpinRera.data());
}
TEST_F(ProgramTests, givenProgramWhenUnknownInternalOptionsArePassedThenTheyAreNotAddedToProgramInternalOptions) {
MockProgram program(toClDeviceVector(*pClDevice));
program.getInternalOptions().erase();
const char *internalOption = "-unknown-internal-options-123";
std::string buildOptions(internalOption);
program.extractInternalOptions(buildOptions);
EXPECT_EQ(0u, program.getInternalOptions().length());
std::string internalOptions;
program.extractInternalOptions(buildOptions, internalOptions);
EXPECT_EQ(0u, internalOptions.length());
}
TEST_F(ProgramTests, givenProgramWhenAllInternalOptionsArePassedMixedWithUnknownInputThenTheyAreParsedCorrectly) {
MockProgram program(toClDeviceVector(*pClDevice));
program.getInternalOptions().erase();
std::string buildOptions = CompilerOptions::concatenate("###", CompilerOptions::gtpinRera, "###", CompilerOptions::greaterThan4gbBuffersRequired, "###");
std::string expectedOutput = CompilerOptions::concatenate(CompilerOptions::gtpinRera, CompilerOptions::greaterThan4gbBuffersRequired);
program.extractInternalOptions(buildOptions);
EXPECT_EQ(expectedOutput, program.getInternalOptions());
std::string internalOptions;
program.extractInternalOptions(buildOptions, internalOptions);
EXPECT_EQ(expectedOutput, internalOptions);
}
TEST_F(ProgramTests, givenProgramWhenInternalOptionsArePassedWithValidValuesThenTheyAreAddedToProgramInternalOptions) {
MockProgram program(toClDeviceVector(*pClDevice));
program.getInternalOptions().erase();
program.isFlagOptionOverride = false;
program.isOptionValueValidOverride = true;
std::string buildOptions = CompilerOptions::concatenate(CompilerOptions::gtpinRera, "someValue");
program.extractInternalOptions(buildOptions);
EXPECT_EQ(buildOptions, program.getInternalOptions()) << program.getInternalOptions();
std::string internalOptions;
program.extractInternalOptions(buildOptions, internalOptions);
EXPECT_EQ(buildOptions, internalOptions) << internalOptions;
}
TEST_F(ProgramTests, givenProgramWhenInternalOptionsArePassedWithInvalidValuesThenTheyAreNotAddedToProgramInternalOptions) {
@ -2577,32 +2553,32 @@ TEST_F(ProgramTests, givenProgramWhenInternalOptionsArePassedWithInvalidValuesTh
program.isFlagOptionOverride = false;
std::string buildOptions = CompilerOptions::concatenate(CompilerOptions::gtpinRera, "someValue");
std::string expectedOutput = "";
program.getInternalOptions().erase();
program.extractInternalOptions(buildOptions);
EXPECT_EQ(expectedOutput, program.getInternalOptions());
std::string internalOptions;
program.extractInternalOptions(buildOptions, internalOptions);
EXPECT_EQ(expectedOutput, internalOptions);
program.isOptionValueValidOverride = true;
buildOptions = std::string(CompilerOptions::gtpinRera);
program.getInternalOptions().erase();
program.extractInternalOptions(buildOptions);
EXPECT_EQ(expectedOutput, program.getInternalOptions());
internalOptions.erase();
program.extractInternalOptions(buildOptions, internalOptions);
EXPECT_EQ(expectedOutput, internalOptions);
}
class AdditionalOptionsMockProgram : public MockProgram {
public:
using MockProgram::MockProgram;
void applyAdditionalOptions() override {
void applyAdditionalOptions(std::string &internalOptions) override {
applyAdditionalOptionsCalled++;
MockProgram::applyAdditionalOptions();
MockProgram::applyAdditionalOptions(internalOptions);
}
uint32_t applyAdditionalOptionsCalled = 0;
};
TEST_F(ProgramTests, givenProgramWhenBuiltThenAdditionalOptionsAreApplied) {
AdditionalOptionsMockProgram program(toClDeviceVector(*pClDevice));
cl_device_id device = pClDevice;
program.build(1, &device, nullptr, false);
program.build(program.getDevices(), nullptr, false);
EXPECT_EQ(1u, program.applyAdditionalOptionsCalled);
}
@ -2776,12 +2752,10 @@ TEST_F(ProgramBinTest, givenPrintProgramBinaryProcessingTimeSetWhenBuildProgramT
DebugManager.flags.PrintProgramBinaryProcessingTime.set(true);
testing::internal::CaptureStdout();
cl_device_id device = pClDevice;
CreateProgramFromBinary(pContext, pContext->getDevices(), "kernel_data_param");
auto retVal = pProgram->build(
1,
&device,
pProgram->getDevices(),
nullptr,
false);
@ -2818,7 +2792,6 @@ struct DebugDataGuard {
TEST_F(ProgramBinTest, GivenBuildWithDebugDataThenBuildDataAvailableViaGetInfo) {
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,
@ -2826,7 +2799,7 @@ TEST_F(ProgramBinTest, GivenBuildWithDebugDataThenBuildDataAvailableViaGetInfo)
&sourceCode,
&knownSourceSize,
retVal);
retVal = pProgram->build(1, &device, nullptr, false);
retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
EXPECT_EQ(CL_SUCCESS, retVal);
// Verify
@ -3016,4 +2989,62 @@ TEST(ProgramCallbackTest, whenInvokeCallbackIsCalledThenFunctionIsProperlyInvoke
EXPECT_TRUE(functionCalled);
program.invokeCallback(nullptr, nullptr);
}
TEST(BuildProgramTest, givenMultiDeviceProgramWhenBuildingThenStoreAndProcessBinaryOnlyOncePerRootDevice) {
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 = clBuildProgram(
pProgram,
0,
nullptr,
nullptr,
nullptr,
nullptr);
for (auto &rootDeviceIndex : context.getRootDeviceIndices()) {
EXPECT_EQ(1, pProgram->replaceDeviceBinaryCalledPerRootDevice[rootDeviceIndex]);
EXPECT_EQ(1, pProgram->processGenBinaryCalledPerRootDevice[rootDeviceIndex]);
}
ASSERT_EQ(CL_SUCCESS, retVal);
retVal = clReleaseProgram(pProgram);
EXPECT_EQ(CL_SUCCESS, retVal);
}

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

@ -57,8 +57,7 @@ TEST_F(ProgramWithBlockKernelsTest, GivenKernelWithBlockKernelsWhenProgramIsBuil
ASSERT_NE(nullptr, mockProgram);
retVal = mockProgram->build(
1,
&device,
pProgram->getDevices(),
nullptr,
false);
EXPECT_EQ(CL_SUCCESS, retVal);

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

@ -56,12 +56,12 @@ TEST(ProgramFromBinary, givenBinaryWithDebugDataWhenCreatingProgramFromBinaryThe
EXPECT_NE(0u, program->getDebugDataSize());
}
class ProgramWithKernelDebuggingTest : public ProgramSimpleFixture,
class ProgramWithKernelDebuggingTest : public ProgramFixture,
public ::testing::Test {
public:
void SetUp() override {
ProgramSimpleFixture::SetUp();
device = pClDevice;
pDevice = static_cast<MockDevice *>(&mockContext.getDevice(0)->getDevice());
if (!pDevice->getHardwareInfo().capabilityTable.debuggerSupported) {
GTEST_SKIP();
}
@ -72,19 +72,17 @@ class ProgramWithKernelDebuggingTest : public ProgramSimpleFixture,
kbHelper = std::make_unique<KernelBinaryHelper>(filename, false);
CreateProgramWithSource(
pContext,
&device,
&mockContext,
"copybuffer.cl");
mockProgram = reinterpret_cast<MockProgram *>(pProgram);
pProgram->enableKernelDebug();
}
void TearDown() override {
ProgramSimpleFixture::TearDown();
ProgramFixture::TearDown();
}
cl_device_id device;
std::unique_ptr<KernelBinaryHelper> kbHelper;
MockProgram *mockProgram = nullptr;
MockUnrestrictiveContext mockContext;
MockDevice *pDevice = nullptr;
};
TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsCompiledThenInternalOptionsIncludeDebugFlag) {
@ -139,7 +137,7 @@ TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsBuilt
debugVars.receivedInternalOptionsOutput = &receivedInternalOptions;
gEnvironment->fclPushDebugVars(debugVars);
cl_int retVal = pProgram->build(1, &device, nullptr, false);
cl_int retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_TRUE(CompilerOptions::contains(receivedInternalOptions, CompilerOptions::debugKernelEnable)) << receivedInternalOptions;
@ -147,7 +145,7 @@ TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsBuilt
}
TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsBuiltThenOptionsIncludeDashGFlag) {
cl_int retVal = pProgram->build(1, &device, nullptr, false);
cl_int retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_THAT(pProgram->getOptions(), ::testing::HasSubstr("-g"));
}
@ -157,7 +155,7 @@ TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugAndOptDisabledWhen
sourceLevelDebugger->isOptDisabled = true;
pDevice->executionEnvironment->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->debugger.reset(sourceLevelDebugger);
cl_int retVal = pProgram->build(1, &device, nullptr, false);
cl_int retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_THAT(pProgram->getOptions(), ::testing::HasSubstr(CompilerOptions::optDisable.data()));
}
@ -167,7 +165,7 @@ TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsBuilt
sourceLevelDebugger->sourceCodeFilename = "debugFileName";
pDevice->executionEnvironment->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->debugger.reset(sourceLevelDebugger);
cl_int retVal = pProgram->build(1, &device, nullptr, false);
cl_int retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_THAT(pProgram->getOptions(), ::testing::StartsWith("-s debugFileName"));
}
@ -179,10 +177,10 @@ TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsLinke
cl_int retVal = pProgram->compile(pProgram->getDevices(), nullptr, 0, nullptr, nullptr);
EXPECT_EQ(CL_SUCCESS, retVal);
auto program = std::unique_ptr<GMockProgram>(new GMockProgram(pContext, false, toClDeviceVector(*pClDevice)));
auto program = std::unique_ptr<GMockProgram>(new GMockProgram(&mockContext, false, mockContext.getDevices()));
program->enableKernelDebug();
EXPECT_CALL(*program, appendKernelDebugOptions()).Times(1);
EXPECT_CALL(*program, appendKernelDebugOptions(::testing::_, ::testing::_)).Times(static_cast<int>(mockContext.getRootDeviceIndices().size()));
cl_program clProgramToLink = pProgram;
retVal = program->link(pProgram->getDevices(), nullptr, 1, &clProgramToLink);
@ -190,32 +188,36 @@ TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsLinke
}
TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsBuiltThenDebuggerIsNotifiedWithKernelDebugData) {
GMockSourceLevelDebugger *sourceLevelDebugger = new GMockSourceLevelDebugger(nullptr);
ON_CALL(*sourceLevelDebugger, notifySourceCode(::testing::_, ::testing::_, ::testing::_)).WillByDefault(::testing::Return(false));
ON_CALL(*sourceLevelDebugger, isOptimizationDisabled()).WillByDefault(::testing::Return(false));
for (auto &rootDeviceIndex : mockContext.getRootDeviceIndices()) {
GMockSourceLevelDebugger *sourceLevelDebugger = new GMockSourceLevelDebugger(nullptr);
ON_CALL(*sourceLevelDebugger, notifySourceCode(::testing::_, ::testing::_, ::testing::_)).WillByDefault(::testing::Return(false));
ON_CALL(*sourceLevelDebugger, isOptimizationDisabled()).WillByDefault(::testing::Return(false));
EXPECT_CALL(*sourceLevelDebugger, isOptimizationDisabled()).Times(1);
EXPECT_CALL(*sourceLevelDebugger, notifySourceCode(::testing::_, ::testing::_, ::testing::_)).Times(1);
EXPECT_CALL(*sourceLevelDebugger, notifyKernelDebugData(::testing::_, ::testing::_, ::testing::_, ::testing::_)).Times(1);
EXPECT_CALL(*sourceLevelDebugger, isOptimizationDisabled()).Times(1);
EXPECT_CALL(*sourceLevelDebugger, notifySourceCode(::testing::_, ::testing::_, ::testing::_)).Times(1);
EXPECT_CALL(*sourceLevelDebugger, notifyKernelDebugData(::testing::_, ::testing::_, ::testing::_, ::testing::_)).Times(1);
sourceLevelDebugger->setActive(true);
pDevice->executionEnvironment->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->debugger.reset(sourceLevelDebugger);
sourceLevelDebugger->setActive(true);
pDevice->executionEnvironment->rootDeviceEnvironments[rootDeviceIndex]->debugger.reset(sourceLevelDebugger);
}
cl_int retVal = pProgram->build(1, &device, nullptr, false);
cl_int retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
EXPECT_EQ(CL_SUCCESS, retVal);
}
TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsLinkedThenDebuggerIsNotifiedWithKernelDebugData) {
GMockSourceLevelDebugger *sourceLevelDebugger = new GMockSourceLevelDebugger(nullptr);
ON_CALL(*sourceLevelDebugger, notifySourceCode(::testing::_, ::testing::_, ::testing::_)).WillByDefault(::testing::Return(false));
ON_CALL(*sourceLevelDebugger, isOptimizationDisabled()).WillByDefault(::testing::Return(false));
for (auto &rootDeviceIndex : mockContext.getRootDeviceIndices()) {
GMockSourceLevelDebugger *sourceLevelDebugger = new GMockSourceLevelDebugger(nullptr);
ON_CALL(*sourceLevelDebugger, notifySourceCode(::testing::_, ::testing::_, ::testing::_)).WillByDefault(::testing::Return(false));
ON_CALL(*sourceLevelDebugger, isOptimizationDisabled()).WillByDefault(::testing::Return(false));
EXPECT_CALL(*sourceLevelDebugger, isOptimizationDisabled()).Times(2);
EXPECT_CALL(*sourceLevelDebugger, notifySourceCode(::testing::_, ::testing::_, ::testing::_)).Times(1);
EXPECT_CALL(*sourceLevelDebugger, notifyKernelDebugData(::testing::_, ::testing::_, ::testing::_, ::testing::_)).Times(1);
EXPECT_CALL(*sourceLevelDebugger, isOptimizationDisabled()).Times(2);
EXPECT_CALL(*sourceLevelDebugger, notifySourceCode(::testing::_, ::testing::_, ::testing::_)).Times(1);
EXPECT_CALL(*sourceLevelDebugger, notifyKernelDebugData(::testing::_, ::testing::_, ::testing::_, ::testing::_)).Times(1);
sourceLevelDebugger->setActive(true);
pDevice->executionEnvironment->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->debugger.reset(sourceLevelDebugger);
sourceLevelDebugger->setActive(true);
pDevice->executionEnvironment->rootDeviceEnvironments[rootDeviceIndex]->debugger.reset(sourceLevelDebugger);
}
cl_int retVal = pProgram->compile(pProgram->getDevices(), nullptr,
0, nullptr, nullptr);
@ -228,7 +230,7 @@ TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsLinke
}
TEST_F(ProgramWithKernelDebuggingTest, givenProgramWithKernelDebugEnabledWhenBuiltThenPatchTokenAllocateSipSurfaceHasSizeGreaterThanZero) {
retVal = pProgram->build(1, &device, CompilerOptions::debugKernelEnable.data(), false);
auto retVal = pProgram->build(pProgram->getDevices(), CompilerOptions::debugKernelEnable.data(), false);
EXPECT_EQ(CL_SUCCESS, retVal);
auto kernelInfo = pProgram->getKernelInfo("CopyBuffer");
@ -236,15 +238,16 @@ TEST_F(ProgramWithKernelDebuggingTest, givenProgramWithKernelDebugEnabledWhenBui
}
TEST_F(ProgramWithKernelDebuggingTest, givenKernelDebugEnabledWhenProgramIsBuiltThenDebugDataIsStored) {
retVal = pProgram->build(1, &device, nullptr, false);
auto retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
EXPECT_EQ(CL_SUCCESS, retVal);
auto debugData = mockProgram->getDebugData();
auto debugData = pProgram->getDebugData();
EXPECT_NE(nullptr, debugData);
EXPECT_NE(0u, mockProgram->getDebugDataSize());
EXPECT_NE(0u, pProgram->getDebugDataSize());
}
TEST_F(ProgramWithKernelDebuggingTest, givenProgramWithKernelDebugEnabledWhenProcessDebugDataIsCalledThenKernelInfosAreFilledWithDebugData) {
retVal = pProgram->build(1, &device, nullptr, false);
auto retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
EXPECT_EQ(CL_SUCCESS, retVal);
pProgram->processDebugData();

1
opencl/test/unit_test/program/program_with_source.h
View File

@ -41,7 +41,6 @@ class ProgramFromSourceTest : public ContextFixture,
CreateProgramWithSource(
pContext,
&device,
SourceFileName);
}