Skip to content

Commit adfbf76

Browse files
jchlandajchlanda
committed
[SYCL] Introduce SYCL_JIT_TARGET_{CPU,FEATURES} env variables
1 parent f1bcc4c commit adfbf76

File tree

8 files changed

+141
-54
lines changed

8 files changed

+141
-54
lines changed

‎sycl-fusion/jit-compiler/include/KernelFusion.h

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -63,10 +63,10 @@ FusionResult fuseKernels(View<SYCLKernelInfo> KernelInformation,
6363
View<ParameterInternalization> Internalization,
6464
View<jit_compiler::JITConstant> JITConstants);
6565

66-
FusionResult
67-
materializeSpecConstants(const char *KernelName,
68-
jit_compiler::SYCLKernelBinaryInfo &BinInfo,
69-
std::vector<unsigned char> &SpecConstBlob);
66+
FusionResult materializeSpecConstants(
67+
const char *KernelName, jit_compiler::SYCLKernelBinaryInfo &BinInfo,
68+
std::vector<unsigned char> &SpecConstBlob, const std::string &TargetCPU,
69+
const std::string &TargetFeatures);
7070

7171
/// Clear all previously set options.
7272
void resetJITConfiguration();

‎sycl-fusion/jit-compiler/lib/KernelFusion.cpp

Lines changed: 6 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -70,10 +70,10 @@ static bool isTargetFormatSupported(BinaryFormat TargetFormat) {
7070
}
7171
}
7272

73-
extern "C" FusionResult
74-
materializeSpecConstants(const char *KernelName,
75-
jit_compiler::SYCLKernelBinaryInfo &BinInfo,
76-
std::vector<unsigned char> &SpecConstBlob) {
73+
extern "C" FusionResult materializeSpecConstants(
74+
const char *KernelName, jit_compiler::SYCLKernelBinaryInfo &BinInfo,
75+
std::vector<unsigned char> &SpecConstBlob, const std::string &TargetCPU,
76+
const std::string &TargetFeatures) {
7777
auto &JITCtx = JITContext::getInstance();
7878

7979
TargetInfo TargetInfo = ConfigHelper::get<option::JITTargetInfo>();
@@ -105,7 +105,8 @@ materializeSpecConstants(const char *KernelName,
105105

106106
SYCLKernelInfo &MaterializerKernelInfo = *ModuleInfo.getKernelFor(KernelName);
107107
if (auto Error = translation::KernelTranslator::translateKernel(
108-
MaterializerKernelInfo, *NewMod, JITCtx, TargetFormat)) {
108+
MaterializerKernelInfo, *NewMod, JITCtx, TargetFormat, TargetCPU,
109+
TargetFeatures)) {
109110
return errorToFusionResult(std::move(Error),
110111
"Translation to output format failed");
111112
}

‎sycl-fusion/jit-compiler/lib/translation/KernelTranslation.cpp

Lines changed: 50 additions & 36 deletions
Original file line numberDiff line numberDiff line change
@@ -168,10 +168,11 @@ KernelTranslator::loadSPIRVKernel(llvm::LLVMContext &LLVMCtx,
168168
return SPIRVLLVMTranslator::loadSPIRVKernel(LLVMCtx, Kernel);
169169
}
170170

171-
llvm::Error KernelTranslator::translateKernel(SYCLKernelInfo &Kernel,
172-
llvm::Module &Mod,
173-
JITContext &JITCtx,
174-
BinaryFormat Format) {
171+
llvm::Error
172+
KernelTranslator::translateKernel(SYCLKernelInfo &Kernel, llvm::Module &Mod,
173+
JITContext &JITCtx, BinaryFormat Format,
174+
const std::string &TargetCPU,
175+
const std::string &TargetFeatures) {
175176

176177
KernelBinary *KernelBin = nullptr;
177178
switch (Format) {
@@ -186,7 +187,7 @@ llvm::Error KernelTranslator::translateKernel(SYCLKernelInfo &Kernel,
186187
}
187188
case BinaryFormat::PTX: {
188189
llvm::Expected<KernelBinary *> BinaryOrError =
189-
translateToPTX(Kernel, Mod, JITCtx);
190+
translateToPTX(Kernel, Mod, JITCtx, TargetCPU, TargetFeatures);
190191
if (auto Error = BinaryOrError.takeError()) {
191192
return Error;
192193
}
@@ -195,7 +196,7 @@ llvm::Error KernelTranslator::translateKernel(SYCLKernelInfo &Kernel,
195196
}
196197
case BinaryFormat::AMDGCN: {
197198
llvm::Expected<KernelBinary *> BinaryOrError =
198-
translateToAMDGCN(Kernel, Mod, JITCtx);
199+
translateToAMDGCN(Kernel, Mod, JITCtx, TargetCPU, TargetFeatures);
199200
if (auto Error = BinaryOrError.takeError())
200201
return Error;
201202
KernelBin = *BinaryOrError;
@@ -226,9 +227,9 @@ KernelTranslator::translateToSPIRV(llvm::Module &Mod, JITContext &JITCtx) {
226227
return SPIRVLLVMTranslator::translateLLVMtoSPIRV(Mod, JITCtx);
227228
}
228229

229-
llvm::Expected<KernelBinary *>
230-
KernelTranslator::translateToPTX(SYCLKernelInfo &KernelInfo, llvm::Module &Mod,
231-
JITContext &JITCtx) {
230+
llvm::Expected<KernelBinary *> KernelTranslator::translateToPTX(
231+
SYCLKernelInfo &KernelInfo, llvm::Module &Mod, JITContext &JITCtx,
232+
const std::string &TargetCPU, const std::string &TargetFeatures) {
232233
#ifndef FUSION_JIT_SUPPORT_PTX
233234
(void)KernelInfo;
234235
(void)Mod;
@@ -257,23 +258,32 @@ KernelTranslator::translateToPTX(SYCLKernelInfo &KernelInfo, llvm::Module &Mod,
257258
ErrorMessage.c_str());
258259
}
259260

260-
llvm::StringRef TargetCPU{"sm_50"};
261-
llvm::StringRef TargetFeatures{"+sm_50,+ptx76"};
262-
if (auto *KernelFunc = Mod.getFunction(KernelInfo.Name.c_str())) {
263-
if (KernelFunc->hasFnAttribute(TARGET_CPU_ATTRIBUTE)) {
264-
TargetCPU =
265-
KernelFunc->getFnAttribute(TARGET_CPU_ATTRIBUTE).getValueAsString();
261+
// Give priority to user specified values (through environment variables:
262+
// SYCL_JIT_TARGET_CPU and SYCL_JIT_TARGET_FEATURES).
263+
llvm::StringRef CPU{TargetCPU};
264+
llvm::StringRef Features{TargetFeatures};
265+
266+
auto *KernelFunc = Mod.getFunction(KernelInfo.Name.c_str());
267+
// If they were not set, use default and consult the module for alternatives
268+
// (if present).
269+
if (CPU.empty()) {
270+
CPU = "sm_50";
271+
if (KernelFunc && KernelFunc->hasFnAttribute(TARGET_CPU_ATTRIBUTE)) {
272+
CPU = KernelFunc->getFnAttribute(TARGET_CPU_ATTRIBUTE).getValueAsString();
266273
}
267-
if (KernelFunc->hasFnAttribute(TARGET_FEATURE_ATTRIBUTE)) {
268-
TargetFeatures = KernelFunc->getFnAttribute(TARGET_FEATURE_ATTRIBUTE)
269-
.getValueAsString();
274+
}
275+
if (Features.empty()) {
276+
Features = "+sm_50,+ptx76";
277+
if (KernelFunc && KernelFunc->hasFnAttribute(TARGET_FEATURE_ATTRIBUTE)) {
278+
Features = KernelFunc->getFnAttribute(TARGET_FEATURE_ATTRIBUTE)
279+
.getValueAsString();
270280
}
271281
}
272282

273283
// FIXME: Check whether we can provide more accurate target information here
274284
auto *TargetMachine = Target->createTargetMachine(
275-
TargetTriple, TargetCPU, TargetFeatures, {}, llvm::Reloc::PIC_,
276-
std::nullopt, llvm::CodeGenOptLevel::Default);
285+
TargetTriple, CPU, Features, {}, llvm::Reloc::PIC_, std::nullopt,
286+
llvm::CodeGenOptLevel::Default);
277287

278288
llvm::legacy::PassManager PM;
279289

@@ -298,9 +308,9 @@ KernelTranslator::translateToPTX(SYCLKernelInfo &KernelInfo, llvm::Module &Mod,
298308
#endif // FUSION_JIT_SUPPORT_PTX
299309
}
300310

301-
llvm::Expected<KernelBinary *>
302-
KernelTranslator::translateToAMDGCN(SYCLKernelInfo &KernelInfo,
303-
llvm::Module &Mod, JITContext &JITCtx) {
311+
llvm::Expected<KernelBinary *> KernelTranslator::translateToAMDGCN(
312+
SYCLKernelInfo &KernelInfo, llvm::Module &Mod, JITContext &JITCtx,
313+
const std::string &TargetCPU, const std::string &TargetFeatures) {
304314
#ifndef FUSION_JIT_SUPPORT_AMDGCN
305315
(void)KernelInfo;
306316
(void)Mod;
@@ -329,25 +339,29 @@ KernelTranslator::translateToAMDGCN(SYCLKernelInfo &KernelInfo,
329339
"Failed to load and translate AMDGCN LLVM IR module with error %s",
330340
ErrorMessage.c_str());
331341

332-
// Set to the lowest tested target according to the GetStartedGuide, section
333-
// "Build DPC++ toolchain with support for HIP AMD"
334-
llvm::StringRef TargetCPU{"gfx906"};
335-
llvm::StringRef TargetFeatures{""};
336-
if (auto *KernelFunc = Mod.getFunction(KernelInfo.Name.c_str())) {
337-
if (KernelFunc->hasFnAttribute(TARGET_CPU_ATTRIBUTE)) {
338-
TargetCPU =
339-
KernelFunc->getFnAttribute(TARGET_CPU_ATTRIBUTE).getValueAsString();
342+
llvm::StringRef CPU{TargetCPU};
343+
llvm::StringRef Features{TargetFeatures};
344+
345+
auto *KernelFunc = Mod.getFunction(KernelInfo.Name.c_str());
346+
if (CPU.empty()) {
347+
// Set to the lowest tested target according to the GetStartedGuide, section
348+
// "Build DPC++ toolchain with support for HIP AMD"
349+
CPU = "gfx906";
350+
if (KernelFunc && KernelFunc->hasFnAttribute(TARGET_CPU_ATTRIBUTE)) {
351+
CPU = KernelFunc->getFnAttribute(TARGET_CPU_ATTRIBUTE).getValueAsString();
340352
}
341-
if (KernelFunc->hasFnAttribute(TARGET_FEATURE_ATTRIBUTE)) {
342-
TargetFeatures = KernelFunc->getFnAttribute(TARGET_FEATURE_ATTRIBUTE)
343-
.getValueAsString();
353+
}
354+
if (Features.empty()) {
355+
if (KernelFunc && KernelFunc->hasFnAttribute(TARGET_FEATURE_ATTRIBUTE)) {
356+
Features = KernelFunc->getFnAttribute(TARGET_FEATURE_ATTRIBUTE)
357+
.getValueAsString();
344358
}
345359
}
346360

347361
// FIXME: Check whether we can provide more accurate target information here
348362
auto *TargetMachine = Target->createTargetMachine(
349-
TargetTriple, TargetCPU, TargetFeatures, {}, llvm::Reloc::PIC_,
350-
std::nullopt, llvm::CodeGenOptLevel::Default);
363+
TargetTriple, CPU, Features, {}, llvm::Reloc::PIC_, std::nullopt,
364+
llvm::CodeGenOptLevel::Default);
351365

352366
std::string AMDObj;
353367
{

‎sycl-fusion/jit-compiler/lib/translation/KernelTranslation.h

Lines changed: 8 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -25,7 +25,9 @@ class KernelTranslator {
2525
loadKernels(llvm::LLVMContext &LLVMCtx, std::vector<SYCLKernelInfo> &Kernels);
2626

2727
static llvm::Error translateKernel(SYCLKernelInfo &Kernel, llvm::Module &Mod,
28-
JITContext &JITCtx, BinaryFormat Format);
28+
JITContext &JITCtx, BinaryFormat Format,
29+
const std::string &TargetCPU = {},
30+
const std::string &TargetFeatures = {});
2931

3032
private:
3133
///
@@ -42,11 +44,14 @@ class KernelTranslator {
4244
JITContext &JITCtx);
4345

4446
static llvm::Expected<KernelBinary *>
45-
translateToPTX(SYCLKernelInfo &Kernel, llvm::Module &Mod, JITContext &JITCtx);
47+
translateToPTX(SYCLKernelInfo &Kernel, llvm::Module &Mod, JITContext &JITCtx,
48+
const std::string &TargetCPU = {},
49+
const std::string &TargetFeatures = {});
4650

4751
static llvm::Expected<KernelBinary *>
4852
translateToAMDGCN(SYCLKernelInfo &KernelInfo, llvm::Module &Mod,
49-
JITContext &JITCtx);
53+
JITContext &JITCtx, const std::string &TargetCPU = {},
54+
const std::string &TargetFeatures = {});
5055
};
5156
} // namespace translation
5257
} // namespace jit_compiler

‎sycl/source/detail/config.def

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -43,3 +43,5 @@ CONFIG(ONEAPI_DEVICE_SELECTOR, 1024, __ONEAPI_DEVICE_SELECTOR)
4343
CONFIG(SYCL_ENABLE_FUSION_CACHING, 1, __SYCL_ENABLE_FUSION_CACHING)
4444
CONFIG(SYCL_CACHE_IN_MEM, 1, __SYCL_CACHE_IN_MEM)
4545
CONFIG(SYCL_JIT_KERNELS, 1, __SYCL_JIT_KERNELS)
46+
CONFIG(SYCL_JIT_TARGET_CPU, 1024, __SYCL_JIT_TARGET_CPU)
47+
CONFIG(SYCL_JIT_TARGET_FEATURES, 1024, __SYCL_JIT_TARGET_FEATURES)

‎sycl/source/detail/config.hpp

Lines changed: 57 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -639,6 +639,63 @@ template <> class SYCLConfig<SYCL_JIT_KERNELS> {
639639
return ValStr;
640640
}
641641
};
642+
643+
template <> class SYCLConfig<SYCL_JIT_TARGET_CPU> {
644+
using BaseT = SYCLConfigBase<SYCL_CACHE_IN_MEM>;
645+
646+
public:
647+
static std::string get() {
648+
const std::string DefaultValue{""};
649+
650+
const char *ValStr = getCachedValue();
651+
652+
if (!ValStr)
653+
return DefaultValue;
654+
655+
return std::string{ValStr};
656+
}
657+
658+
static void reset() { (void)getCachedValue(/*ResetCache=*/true); }
659+
660+
static const char *getName() { return BaseT::MConfigName; }
661+
662+
private:
663+
static const char *getCachedValue(bool ResetCache = false) {
664+
static const char *ValStr = BaseT::getRawValue();
665+
if (ResetCache)
666+
ValStr = BaseT::getRawValue();
667+
return ValStr;
668+
}
669+
};
670+
671+
template <> class SYCLConfig<SYCL_JIT_TARGET_FEATURES> {
672+
using BaseT = SYCLConfigBase<SYCL_CACHE_IN_MEM>;
673+
674+
public:
675+
static std::string get() {
676+
const std::string DefaultValue{""};
677+
678+
const char *ValStr = getCachedValue();
679+
680+
if (!ValStr)
681+
return DefaultValue;
682+
683+
return std::string{ValStr};
684+
}
685+
686+
static void reset() { (void)getCachedValue(/*ResetCache=*/true); }
687+
688+
static const char *getName() { return BaseT::MConfigName; }
689+
690+
private:
691+
static const char *getCachedValue(bool ResetCache = false) {
692+
static const char *ValStr = BaseT::getRawValue();
693+
if (ResetCache)
694+
ValStr = BaseT::getRawValue();
695+
return ValStr;
696+
}
697+
};
698+
642699
#undef INVALID_CONFIG_EXCEPTION
643700

644701
} // namespace detail

‎sycl/source/detail/jit_compiler.cpp

Lines changed: 7 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -726,7 +726,6 @@ sycl::detail::pi::PiKernel jit_compiler::materializeSpecConstants(
726726
BinaryImageFormat, 0, RawDeviceImage.BinaryStart, DeviceImageSize};
727727

728728
::jit_compiler::TargetInfo TargetInfo = getTargetInfo(Queue);
729-
::jit_compiler::BinaryFormat TargetFormat = TargetInfo.getFormat();
730729
AddToConfigHandle(
731730
::jit_compiler::option::JITTargetInfo::set(std::move(TargetInfo)));
732731
bool DebugEnabled =
@@ -736,8 +735,13 @@ sycl::detail::pi::PiKernel jit_compiler::materializeSpecConstants(
736735
AddToConfigHandle(::jit_compiler::option::JITEnableCaching::set(
737736
detail::SYCLConfig<detail::SYCL_ENABLE_FUSION_CACHING>::get()));
738737

739-
auto MaterializerResult =
740-
MaterializeSpecConstHandle(KernelName.c_str(), BinInfo, SpecConstBlob);
738+
std::string TargetCPU =
739+
detail::SYCLConfig<detail::SYCL_JIT_TARGET_CPU>::get();
740+
std::string TargetFeatures =
741+
detail::SYCLConfig<detail::SYCL_JIT_TARGET_FEATURES>::get();
742+
743+
auto MaterializerResult = MaterializeSpecConstHandle(
744+
KernelName.c_str(), BinInfo, SpecConstBlob, TargetCPU, TargetFeatures);
741745
if (MaterializerResult.failed()) {
742746
std::string Message{"Compilation for kernel failed with message:\n"};
743747
Message.append(MaterializerResult.getErrorMessage());

‎sycl/source/detail/program_manager/program_manager.cpp

Lines changed: 7 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -2454,10 +2454,14 @@ sycl::detail::pi::PiKernel ProgramManager::getOrCreateMaterializedKernel(
24542454
auto &Plugin = DeviceImpl->getPlugin();
24552455
ProgramPtr ProgramManaged(
24562456
Program, Plugin->getPiPlugin().PiFunctionTable.piProgramRelease);
2457-
// TODO: JKB: Flags and zeros.
2457+
2458+
std::string CompileOpts;
2459+
std::string LinkOpts;
2460+
applyOptionsFromEnvironment(CompileOpts, LinkOpts);
24582461
auto BuildProgram =
2459-
build(std::move(ProgramManaged), detail::getSyclObjImpl(Context), "", "",
2460-
DeviceImpl->getHandleRef(), 0);
2462+
build(std::move(ProgramManaged), detail::getSyclObjImpl(Context),
2463+
CompileOpts, LinkOpts, DeviceImpl->getHandleRef(),
2464+
/*For non SPIR-V devices DeviceLibReqdMask is always 0*/ 0);
24612465
sycl::detail::pi::PiKernel PiKernel{nullptr};
24622466
Plugin->call<errc::kernel_not_supported, PiApiKind::piKernelCreate>(
24632467
BuildProgram.get(), KernelName.c_str(), &PiKernel);

0 commit comments

Comments
 (0)