Skip to content

[SYCL][RTC] Rework handling of build_options #17405

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 6 commits into from
Mar 19, 2025
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
19 changes: 12 additions & 7 deletions sycl-jit/jit-compiler/include/KernelFusion.h
Original file line number Diff line number Diff line change
Expand Up @@ -86,32 +86,37 @@ class RTCHashResult {
sycl::detail::string HashOrLog;
};

enum class RTCErrorCode { SUCCESS, BUILD, INVALID };

class RTCResult {
public:
explicit RTCResult(const char *BuildLog)
: Failed{true}, BundleInfo{}, BuildLog{BuildLog} {}
explicit RTCResult(const char *BuildLog,
RTCErrorCode ErrorCode = RTCErrorCode::BUILD)
: ErrorCode{ErrorCode}, BundleInfo{}, BuildLog{BuildLog} {
assert(ErrorCode != RTCErrorCode::SUCCESS);
}

RTCResult(RTCBundleInfo &&BundleInfo, RTCDeviceCodeIR &&DeviceCodeIR,
const char *BuildLog)
: Failed{false}, BundleInfo{std::move(BundleInfo)},
: ErrorCode{RTCErrorCode::SUCCESS}, BundleInfo{std::move(BundleInfo)},
DeviceCodeIR(std::move(DeviceCodeIR)), BuildLog{BuildLog} {}

bool failed() const { return Failed; }
RTCErrorCode getErrorCode() const { return ErrorCode; }

const char *getBuildLog() const { return BuildLog.c_str(); }

const RTCBundleInfo &getBundleInfo() const {
assert(!failed() && "No bundle info");
assert(ErrorCode == RTCErrorCode::SUCCESS && "No bundle info");
return BundleInfo;
}

const RTCDeviceCodeIR &getDeviceCodeIR() const {
assert(!failed() && "No device code IR");
assert(ErrorCode == RTCErrorCode::SUCCESS && "No device code IR");
return DeviceCodeIR;
}

private:
bool Failed;
RTCErrorCode ErrorCode;
RTCBundleInfo BundleInfo;
RTCDeviceCodeIR DeviceCodeIR;
sycl::detail::string BuildLog;
Expand Down
10 changes: 6 additions & 4 deletions sycl-jit/jit-compiler/lib/KernelFusion.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -45,11 +45,12 @@ static std::string formatError(llvm::Error &&Err, const std::string &Msg) {
return ErrMsg.str();
}

template <typename ResultType>
static ResultType errorTo(llvm::Error &&Err, const std::string &Msg) {
template <typename ResultType, typename... ExtraArgTypes>
static ResultType errorTo(llvm::Error &&Err, const std::string &Msg,
ExtraArgTypes... ExtraArgs) {
// Cannot throw an exception here if LLVM itself is compiled without exception
// support.
return ResultType{formatError(std::move(Err), Msg).c_str()};
return ResultType{formatError(std::move(Err), Msg).c_str(), ExtraArgs...};
}

static std::vector<jit_compiler::NDRange>
Expand Down Expand Up @@ -288,7 +289,8 @@ compileSYCL(InMemoryFile SourceFile, View<InMemoryFile> IncludeFiles,
auto UserArgListOrErr = parseUserArgs(UserArgs);
if (!UserArgListOrErr) {
return errorTo<RTCResult>(UserArgListOrErr.takeError(),
"Parsing of user arguments failed");
"Parsing of user arguments failed",
RTCErrorCode::INVALID);
}
llvm::opt::InputArgList UserArgList = std::move(*UserArgListOrErr);

Expand Down
95 changes: 27 additions & 68 deletions sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -321,20 +321,15 @@ static void adjustArgs(const InputArgList &UserArgList,
DAL.AddJoinedArg(
nullptr, OptTable.getOption(OPT_resource_dir_EQ),
(DPCPPRoot + "/lib/clang/" + Twine(CLANG_VERSION_MAJOR)).str());
for (auto *Arg : UserArgList) {
DAL.append(Arg);
}
// Remove args that will trigger an unused command line argument warning for
// the FrontendAction invocation, but are handled later (e.g. during device
// linking).
DAL.eraseArg(OPT_fsycl_device_lib_EQ);
DAL.eraseArg(OPT_fno_sycl_device_lib_EQ);
DAL.eraseArg(OPT_ftime_trace_EQ);
DAL.eraseArg(OPT_ftime_trace_granularity_EQ);
DAL.eraseArg(OPT_ftime_trace_verbose);
// User args may contain options not intended for the frontend, but we can't
// claim them here to tell the driver they're used later. Hence, suppress the
// unused argument warning.
DAL.AddFlagArg(nullptr, OptTable.getOption(OPT_Qunused_arguments));

ArgStringList ASL;
for_each(DAL, [&DAL, &ASL](Arg *A) { A->render(DAL, ASL); });
for_each(UserArgList,
[&UserArgList, &ASL](Arg *A) { A->render(UserArgList, ASL); });
transform(ASL, std::back_inserter(CommandLine),
[](const char *AS) { return std::string{AS}; });
}
Expand Down Expand Up @@ -751,51 +746,21 @@ jit_compiler::parseUserArgs(View<const char *> UserArgs) {
UserArgsRef[MissingArgIndex], MissingArgIndex);
}

// Check for unsupported options.
// TODO: There are probably more, e.g. requesting non-SPIR-V targets.
{
// -fsanitize=address
bool IsDeviceAsanEnabled = false;
if (Arg *A = AL.getLastArg(OPT_fsanitize_EQ, OPT_fno_sanitize_EQ)) {
if (A->getOption().matches(OPT_fsanitize_EQ) &&
A->getValues().size() == 1) {
std::string SanitizeVal = A->getValue();
IsDeviceAsanEnabled = SanitizeVal == "address";
}
} else {
// User can pass -fsanitize=address to device compiler via
// -Xsycl-target-frontend.
auto SyclFEArg = AL.getAllArgValues(OPT_Xsycl_frontend);
IsDeviceAsanEnabled = (std::count(SyclFEArg.begin(), SyclFEArg.end(),
"-fsanitize=address") > 0);
if (!IsDeviceAsanEnabled) {
auto SyclFEArgEq = AL.getAllArgValues(OPT_Xsycl_frontend_EQ);
IsDeviceAsanEnabled =
(std::count(SyclFEArgEq.begin(), SyclFEArgEq.end(),
"-fsanitize=address") > 0);
}

// User can also enable asan for SYCL device via -Xarch_device option.
if (!IsDeviceAsanEnabled) {
auto DeviceArchVals = AL.getAllArgValues(OPT_Xarch_device);
for (auto DArchVal : DeviceArchVals) {
if (DArchVal.find("-fsanitize=address") != std::string::npos) {
IsDeviceAsanEnabled = true;
break;
}
}
}
}

if (IsDeviceAsanEnabled) {
return createStringError(
"Device ASAN is not supported for runtime compilation");
}
}

if (!AL.hasFlag(OPT_fsycl_device_code_split_esimd,
OPT_fno_sycl_device_code_split_esimd, true)) {
return createStringError("ESIMD device code split cannot be deactivated");
// Check for options that are unsupported because they would interfere with
// the in-memory pipeline.
Arg *UnsupportedArg =
AL.getLastArg(OPT_Action_Group, // Actions like -c or -S
OPT_Link_Group, // Linker flags
OPT_o, // Output file
OPT_fsycl_targets_EQ, // AoT compilation
OPT_fsycl_link_EQ, // SYCL linker
OPT_fno_sycl_device_code_split_esimd, // invoke_simd
OPT_fsanitize_EQ // Sanitizer
);
if (UnsupportedArg) {
return createStringError(
"Option '%s' is not supported for SYCL runtime compilation",
UnsupportedArg->getAsString(AL).c_str());
}

return std::move(AL);
Expand All @@ -806,20 +771,14 @@ void jit_compiler::encodeBuildOptions(RTCBundleInfo &BundleInfo,
std::string CompileOptions;
raw_string_ostream COSOS{CompileOptions};

for (Arg *A : UserArgList.getArgs()) {
if (!(A->getOption().matches(OPT_Xs) ||
A->getOption().matches(OPT_Xs_separate))) {
continue;
for (Arg *A : UserArgList.filtered(OPT_Xs, OPT_Xs_separate)) {
if (!CompileOptions.empty()) {
COSOS << ' ';
}

// Trim first and last quote if they exist, but no others.
StringRef AV{A->getValue()};
AV = AV.trim();
if (AV.front() == AV.back() && (AV.front() == '\'' || AV.front() == '"')) {
AV = AV.drop_front().drop_back();
if (A->getOption().matches(OPT_Xs)) {
COSOS << '-';
}

COSOS << (CompileOptions.empty() ? "" : " ") << AV;
COSOS << A->getValue();
}

if (!CompileOptions.empty()) {
Expand Down
13 changes: 10 additions & 3 deletions sycl/source/detail/jit_compiler.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -1302,9 +1302,16 @@ std::pair<sycl_device_binaries, std::string> jit_compiler::compileSYCL(
auto Result = CompileSYCLHandle(SourceFile, IncludeFilesView, UserArgsView,
CachedIR, /*SaveIR=*/!CacheKey.empty());

appendToLog(Result.getBuildLog());
if (Result.failed()) {
throw sycl::exception(sycl::errc::build, Result.getBuildLog());
const char *BuildLog = Result.getBuildLog();
appendToLog(BuildLog);
switch (Result.getErrorCode()) {
using RTCErrC = ::jit_compiler::RTCErrorCode;
case RTCErrC::BUILD:
throw sycl::exception(sycl::errc::build, BuildLog);
case RTCErrC::INVALID:
throw sycl::exception(sycl::errc::invalid, BuildLog);
default: // RTCErrC::SUCCESS
break;
}

const auto &IR = Result.getDeviceCodeIR();
Expand Down
28 changes: 11 additions & 17 deletions sycl/test-e2e/KernelCompiler/sycl.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -551,23 +551,17 @@ int test_unsupported_options() {
ctx, syclex::source_language::sycl, "");
std::vector<sycl::device> devs = kbSrc.get_devices();

auto CheckUnsupported = [&](const std::vector<std::string> &flags) {
try {
syclex::build(kbSrc, devs,
syclex::properties{syclex::build_options{flags}});
assert(false && "unsupported option not detected");
} catch (sycl::exception &e) {
assert(e.code() == sycl::errc::build);
assert(std::string(e.what()).find("Parsing of user arguments failed") !=
std::string::npos);
}
};

CheckUnsupported({"-fsanitize=address"});
CheckUnsupported({"-Xsycl-target-frontend", "-fsanitize=address"});
CheckUnsupported({"-Xsycl-target-frontend=spir64", "-fsanitize=address"});
CheckUnsupported({"-Xarch_device", "-fsanitize=address"});
CheckUnsupported({"-fno-sycl-device-code-split-esimd"});
try {
// Don't attempt to test exhaustively here...
syclex::build(kbSrc, devs,
syclex::properties{
syclex::build_options{"-fsycl-targets=intel_gpu_pvc"}});
assert(false && "unsupported option not detected");
} catch (sycl::exception &e) {
assert(e.code() == sycl::errc::invalid);
assert(std::string(e.what()).find("Parsing of user arguments failed") !=
std::string::npos);
}

return 0;
}
Expand Down
5 changes: 2 additions & 3 deletions sycl/test-e2e/KernelCompiler/sycl_device_flags.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -129,9 +129,8 @@ int main(int argc, char *argv[]) {
source_kb kbSrc = syclex::create_kernel_bundle_from_source(
ctx, syclex::source_language::sycl, SYCLSource);

// Flags with and without space, inner quotes.
std::vector<std::string> flags{"-Xs '-doubleGRF'",
"-Xs'-Xfinalizer \"-printregusage\"'"};
std::vector<std::string> flags{"-Xs", "-doubleGRF",
"-XsXfinalizer \"-printregusage\""};
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@cperkinsintel This occurred to me while working on the documentation:

  • Do we still need to support the implicit unquoting, given build_options are always a vector? I didn't find that in the "normal" option handling (but might have overlooked it).
  • The -Xs option says that for the joined form, a dash is inserted automatically, so -XsDfoo becomes -Dfoo when passed to the target compiler.

exe_kb kbExe =
syclex::build(kbSrc, syclex::properties{syclex::build_options{flags}});

Expand Down