Skip to content

Commit

Permalink
Merge remote-tracking branch 'upstream/sycl' into dump_record_layout
Browse files Browse the repository at this point in the history 
* upstream/sycl:
  [SYCL] Handle KernelName templated using type with enum template argument (intel#1780)
  [SYCL] Fix KernelNameInfo generated for empty template parameter pack (intel#1775)
  [SYCL] Do not export utility methods from SYCLMemObjT (intel#1768)
  [Driver][SYCL] Fix processing when using -fsycl-link (intel#1765)
  [SYCL][NFC] Remove outdated confusing comment (intel#1779)
  [SYCL][NFC] Wrap classes in .cpp into a namespace to disable external linkage. (intel#1776)
  [SYCL][CUDA] Fixes CUDA unit tests that uses SYCL directly (intel#1763)
  [SYCL][Doc] Fix default device selection rules doc (intel#1769)
  [SYCL][CUDA] Remove pi Event Callback implementation (intel#1735)
  [SYCL] Throw exception if range/offset of kernel execution exceeds INT_MAX (intel#1713)
  [SYCL-PTX] Add intermediate layer to libclc to ease type management (intel#1712)
  • Loading branch information
Alexander Batashev committed Jun 2, 2020
2 parents e2c9419 + f9226d2 commit fc3b63d
Show file tree
Hide file tree
Showing 169 changed files with 15,424 additions and 2,306 deletions.
7 changes: 5 additions & 2 deletions clang/include/clang/Sema/Sema.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -319,7 +319,8 @@ class SYCLIntegrationHeader {
};

public:
SYCLIntegrationHeader(DiagnosticsEngine &Diag, bool UnnamedLambdaSupport);
SYCLIntegrationHeader(DiagnosticsEngine &Diag, bool UnnamedLambdaSupport,
Sema &S);

/// Emits contents of the header into given stream.
void emit(raw_ostream &Out);
Expand Down Expand Up @@ -424,6 +425,8 @@ class SYCLIntegrationHeader {

/// Whether header is generated with unnamed lambda support
bool UnnamedLambdaSupport;

Sema &S;
};

/// Keeps track of expected type during expression parsing. The type is tied to
Expand Down Expand Up @@ -12584,7 +12587,7 @@ class Sema final {
SYCLIntegrationHeader &getSyclIntegrationHeader() {
if (SyclIntHeader == nullptr)
SyclIntHeader = std::make_unique<SYCLIntegrationHeader>(
getDiagnostics(), getLangOpts().SYCLUnnamedLambda);
getDiagnostics(), getLangOpts().SYCLUnnamedLambda, *this);
return *SyclIntHeader.get();
}

Expand Down
4 changes: 4 additions & 0 deletions clang/lib/CodeGen/CGSYCLRuntime.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -20,6 +20,8 @@
using namespace clang;
using namespace CodeGen;

namespace {

/// Various utilities.
/// TODO partially duplicates functionality from SemaSYCL.cpp, can be shared.
class Util {
Expand Down Expand Up @@ -52,6 +54,8 @@ const char *WG_SCOPE_MD_ID = "work_group_scope";
const char *WI_SCOPE_MD_ID = "work_item_scope";
const char *PFWI_MD_ID = "parallel_for_work_item";

} // anonymous namespace

bool CGSYCLRuntime::actOnFunctionStart(const FunctionDecl &FD,
llvm::Function &F) {
SYCLScopeAttr *Scope = FD.getAttr<SYCLScopeAttr>();
Expand Down
21 changes: 15 additions & 6 deletions clang/lib/Driver/Driver.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -3549,12 +3549,21 @@ class OffloadingActionBuilder final {
for (auto SDA : SYCLDeviceActions)
SYCLLinkBinaryList.push_back(SDA);
if (WrapDeviceOnlyBinary) {
auto *DeviceLinkAction =
C.MakeAction<LinkJobAction>(SYCLLinkBinaryList, types::TY_Image);
// Wrap the binary when -fsycl-link is given
SYCLLinkBinary =
C.MakeAction<OffloadWrapperJobAction>(DeviceLinkAction,
types::TY_Object);
// -fsycl-link behavior does the following to the unbundled device
// binaries:
// 1) Link them together using llvm-link
// 2) Pass the linked binary through sycl-post-link
// 3) Translate final .bc file to .spv
// 4) Wrap the binary with the offload wrapper which can be used
// by any compilation link step.
auto *DeviceLinkAction = C.MakeAction<LinkJobAction>(
SYCLLinkBinaryList, types::TY_Image);
auto *PostLinkAction = C.MakeAction<SYCLPostLinkJobAction>(
DeviceLinkAction, types::TY_LLVM_BC);
auto *TranslateAction = C.MakeAction<SPIRVTranslatorJobAction>(
PostLinkAction, types::TY_Image);
SYCLLinkBinary = C.MakeAction<OffloadWrapperJobAction>(
TranslateAction, types::TY_Object);
} else {
auto *Link = C.MakeAction<LinkJobAction>(SYCLLinkBinaryList,
types::TY_Image);
Expand Down
42 changes: 27 additions & 15 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -54,6 +54,8 @@ enum KernelInvocationKind {
const static std::string InitMethodName = "__init";
const static std::string FinalizeMethodName = "__finalize";

namespace {

/// Various utilities.
class Util {
public:
Expand Down Expand Up @@ -91,6 +93,8 @@ class Util {
ArrayRef<Util::DeclContextDesc> Scopes);
};

} // anonymous namespace

// This information is from Section 4.13 of the SYCL spec
// https://www.khronos.org/registry/SYCL/specs/sycl-1.2.1.pdf
// This function returns false if the math lib function
Expand Down Expand Up @@ -1454,7 +1458,6 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc,
}

void Sema::MarkDevice(void) {
// Let's mark all called functions with SYCL Device attribute.
// Create the call graph so we can detect recursion and check the validity
// of new operator overrides. Add the kernel function itself in case
// it is recursive.
Expand Down Expand Up @@ -1863,6 +1866,9 @@ static void printArguments(ASTContext &Ctx, raw_ostream &ArgOS,
ArrayRef<TemplateArgument> Args,
const PrintingPolicy &P);

static std::string getKernelNameTypeString(QualType T, ASTContext &Ctx,
const PrintingPolicy &TypePolicy);

static void printArgument(ASTContext &Ctx, raw_ostream &ArgOS,
TemplateArgument Arg, const PrintingPolicy &P) {
switch (Arg.getKind()) {
Expand All @@ -1888,8 +1894,7 @@ static void printArgument(ASTContext &Ctx, raw_ostream &ArgOS,
TypePolicy.SuppressTypedefs = true;
TypePolicy.SuppressTagKeyword = true;
QualType T = Arg.getAsType();
QualType FullyQualifiedType = TypeName::getFullyQualifiedType(T, Ctx, true);
ArgOS << FullyQualifiedType.getAsString(TypePolicy);
ArgOS << getKernelNameTypeString(T, Ctx, TypePolicy);
break;
}
default:
Expand All @@ -1903,6 +1908,10 @@ static void printArguments(ASTContext &Ctx, raw_ostream &ArgOS,
for (unsigned I = 0; I < Args.size(); I++) {
const TemplateArgument &Arg = Args[I];

// If argument is an empty pack argument, skip printing comma and argument.
if (Arg.getKind() == TemplateArgument::ArgKind::Pack && !Arg.pack_size())
continue;

if (I != 0)
ArgOS << ", ";

Expand All @@ -1918,36 +1927,36 @@ static void printTemplateArguments(ASTContext &Ctx, raw_ostream &ArgOS,
ArgOS << ">";
}

static std::string getKernelNameTypeString(QualType T) {
static std::string getKernelNameTypeString(QualType T, ASTContext &Ctx,
const PrintingPolicy &TypePolicy) {

QualType FullyQualifiedType = TypeName::getFullyQualifiedType(T, Ctx, true);

const CXXRecordDecl *RD = T->getAsCXXRecordDecl();

if (!RD)
return getCPPTypeString(T);
return eraseAnonNamespace(FullyQualifiedType.getAsString(TypePolicy));

// If kernel name type is a template specialization with enum type
// template parameters, enumerators in name type string should be
// replaced with their underlying value since the enum definition
// is not visible in integration header.
if (const auto *TSD = dyn_cast<ClassTemplateSpecializationDecl>(RD)) {
LangOptions LO;
PrintingPolicy P(LO);
P.SuppressTypedefs = true;
SmallString<64> Buf;
llvm::raw_svector_ostream ArgOS(Buf);

// Print template class name
TSD->printQualifiedName(ArgOS, P, /*WithGlobalNsPrefix*/ true);
TSD->printQualifiedName(ArgOS, TypePolicy, /*WithGlobalNsPrefix*/ true);

// Print template arguments substituting enumerators
ASTContext &Ctx = RD->getASTContext();
const TemplateArgumentList &Args = TSD->getTemplateArgs();
printTemplateArguments(Ctx, ArgOS, Args.asArray(), P);
printTemplateArguments(Ctx, ArgOS, Args.asArray(), TypePolicy);

return eraseAnonNamespace(ArgOS.str().str());
}

return getCPPTypeString(T);
return eraseAnonNamespace(FullyQualifiedType.getAsString(TypePolicy));
}

void SYCLIntegrationHeader::emit(raw_ostream &O) {
Expand Down Expand Up @@ -2066,9 +2075,11 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
O << "', '" << c;
O << "'> {\n";
} else {

LangOptions LO;
PrintingPolicy P(LO);
P.SuppressTypedefs = true;
O << "template <> struct KernelInfo<"
<< getKernelNameTypeString(K.NameType) << "> {\n";
<< getKernelNameTypeString(K.NameType, S.getASTContext(), P) << "> {\n";
}
O << " DLL_LOCAL\n";
O << " static constexpr const char* getName() { return \"" << K.Name
Expand Down Expand Up @@ -2137,8 +2148,9 @@ void SYCLIntegrationHeader::addSpecConstant(StringRef IDName, QualType IDType) {
}

SYCLIntegrationHeader::SYCLIntegrationHeader(DiagnosticsEngine &_Diag,
bool _UnnamedLambdaSupport)
: Diag(_Diag), UnnamedLambdaSupport(_UnnamedLambdaSupport) {}
bool _UnnamedLambdaSupport,
Sema &_S)
: Diag(_Diag), UnnamedLambdaSupport(_UnnamedLambdaSupport), S(_S) {}

// -----------------------------------------------------------------------------
// Utility class methods
Expand Down
9 changes: 9 additions & 0 deletions clang/test/CodeGenSYCL/int_header1.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -12,6 +12,7 @@
// CHECK:template <> struct KernelInfo<::nm1::KernelName4<KernelName7>> {
// CHECK:template <> struct KernelInfo<::nm1::KernelName8<::nm1::nm2::C>> {
// CHECK:template <> struct KernelInfo<::TmplClassInAnonNS<ClassInAnonNS>> {
// CHECK:template <> struct KernelInfo<::nm1::KernelName9<char>> {

// This test checks if the SYCL device compiler is able to generate correct
// integration header when the kernel name class is expressed in different
Expand Down Expand Up @@ -42,6 +43,9 @@ namespace nm1 {
template <> class KernelName4<nm1::nm2::KernelName0> {};
template <> class KernelName4<KernelName1> {};

template <typename T, typename...>
class KernelName9;

} // namespace nm1

namespace {
Expand Down Expand Up @@ -128,6 +132,10 @@ struct MyWrapper {
kernel_single_task<TmplClassInAnonNS<class ClassInAnonNS>>(
[=]() { acc.use(); });

// Kernel name type is a templated specialization class with empty template pack argument
kernel_single_task<nm1::KernelName9<char>>(
[=]() { acc.use(); });

return 0;
}
};
Expand All @@ -151,5 +159,6 @@ int main() {
KernelInfo<class nm1::KernelName4<class KernelName7>>::getName();
KernelInfo<class nm1::KernelName8<nm1::nm2::C>>::getName();
KernelInfo<class TmplClassInAnonNS<class ClassInAnonNS>>::getName();
KernelInfo<class nm1::KernelName9<char>>::getName();
#endif //__SYCL_DEVICE_ONLY__
}
31 changes: 30 additions & 1 deletion clang/test/CodeGenSYCL/kernelname-enum.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -79,6 +79,21 @@ class dummy_functor_7 {
void operator()() {}
};

namespace type_argument_template_enum {
enum class E : int {
A,
B,
C
};
}

template <typename T>
class T1 {};
template <type_argument_template_enum::E EnumValue>
class T2 {};
template <typename EnumType>
class T3 {};

int main() {

dummy_functor_1<no_namespace_int::val_1> f1;
Expand Down Expand Up @@ -124,6 +139,14 @@ int main() {
cgh.single_task(f8);
});

q.submit([&](cl::sycl::handler &cgh) {
cgh.single_task<T1<T2<type_argument_template_enum::E::A>>>([=]() {});
});

q.submit([&](cl::sycl::handler &cgh) {
cgh.single_task<T1<T3<type_argument_template_enum::E>>>([=]() {});
});

return 0;
}

Expand All @@ -145,7 +168,11 @@ int main() {
// CHECK: enum unscoped_enum : int;
// CHECK: template <unscoped_enum EnumType> class dummy_functor_6;
// CHECK: template <typename EnumType> class dummy_functor_7;

// CHECK: namespace type_argument_template_enum {
// CHECK-NEXT: enum class E : int;
// CHECK-NEXT: }
// CHECK: template <type_argument_template_enum::E EnumValue> class T2;
// CHECK: template <typename T> class T1;
// CHECK: Specializations of KernelInfo for kernel function types:
// CHECK: template <> struct KernelInfo<::dummy_functor_1<(no_namespace_int)0>>
// CHECK: template <> struct KernelInfo<::dummy_functor_2<(no_namespace_short)1>>
Expand All @@ -155,3 +182,5 @@ int main() {
// CHECK: template <> struct KernelInfo<::dummy_functor_6<(unscoped_enum)0>>
// CHECK: template <> struct KernelInfo<::dummy_functor_7<::no_namespace_int>>
// CHECK: template <> struct KernelInfo<::dummy_functor_7<::internal::namespace_short>>
// CHECK: template <> struct KernelInfo<::T1<::T2<(type_argument_template_enum::E)0>>>
// CHECK: template <> struct KernelInfo<::T1<::T3<::type_argument_template_enum::E>>>
12 changes: 8 additions & 4 deletions clang/test/Driver/sycl-offload.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -423,8 +423,10 @@
// CHK-LINK-UB: 0: input, "[[INPUT:.+\.o]]", object
// CHK-LINK-UB: 1: clang-offload-unbundler, {0}, object
// CHK-LINK-UB: 2: linker, {1}, image, (device-sycl)
// CHK-LINK-UB: 3: clang-offload-wrapper, {2}, object, (device-sycl)
// CHK-LINK-UB: 4: offload, "device-sycl (spir64-unknown-unknown-sycldevice{{.*}})" {3}, object
// CHK-LINK-UB: 3: sycl-post-link, {2}, ir, (device-sycl)
// CHK-LINK-UB: 4: llvm-spirv, {3}, image, (device-sycl)
// CHK-LINK-UB: 5: clang-offload-wrapper, {4}, object, (device-sycl)
// CHK-LINK-UB: 6: offload, "device-sycl (spir64-unknown-unknown-sycldevice)" {5}, object

/// ###########################################################################

Expand All @@ -437,8 +439,10 @@
// CHK-LINK: 1: preprocessor, {0}, cpp-output, (device-sycl)
// CHK-LINK: 2: compiler, {1}, ir, (device-sycl)
// CHK-LINK: 3: linker, {2}, image, (device-sycl)
// CHK-LINK: 4: clang-offload-wrapper, {3}, object, (device-sycl)
// CHK-LINK: 5: offload, "device-sycl (spir64-unknown-unknown-sycldevice{{.*}})" {4}, object
// CHK-LINK: 4: sycl-post-link, {3}, ir, (device-sycl)
// CHK-LINK: 5: llvm-spirv, {4}, image, (device-sycl)
// CHK-LINK: 6: clang-offload-wrapper, {5}, object, (device-sycl)
// CHK-LINK: 7: offload, "device-sycl (spir64-unknown-unknown-sycldevice)" {6}, object

/// ###########################################################################

Expand Down
13 changes: 10 additions & 3 deletions libclc/CMakeLists.txt
View file
Original file line number Diff line number Diff line change
Expand Up @@ -160,8 +160,15 @@ endif()

find_program( PYTHON python )
file( TO_CMAKE_PATH ${CMAKE_CURRENT_SOURCE_DIR}/generic/lib/gen_convert.py clc_script_loc )
file( TO_CMAKE_PATH ${CMAKE_CURRENT_SOURCE_DIR}/generic/libspirv/gen_core_convert.py core_script_loc )
file( TO_CMAKE_PATH ${CMAKE_CURRENT_SOURCE_DIR}/generic/libspirv/gen_convert.py spirv_script_loc )

add_custom_command(
OUTPUT convert-core.cl
COMMAND ${PYTHON} ${core_script_loc} > convert-core.cl
DEPENDS ${core_script_loc} )
add_custom_target( "generate_convert_core.cl" DEPENDS convert-core.cl )

add_custom_command(
OUTPUT convert-spirv.cl
COMMAND ${PYTHON} ${spirv_script_loc} > convert-spirv.cl
Expand Down Expand Up @@ -211,7 +218,7 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
libclc_configure_lib_source(libspirv_files
LIB_DIR libspirv
DIRS ${dirs} ${DARCH} ${DARCH}-${OS} ${DARCH}-${VENDOR}-${OS}
DEPS convert-spirv.cl )
DEPS convert-spirv.cl convert-core.cl)

foreach( d ${${t}_devices} )
# Some targets don't have a specific GPU to target
Expand All @@ -230,7 +237,7 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
COMPILE_OPT ${mcpu}
FILES ${libspirv_files}
ALIASES ${${d}_aliases}
GENERATE_TARGET "generate_convert_clc.cl"
GENERATE_TARGET "generate_convert_spirv.cl" "generate_convert_core.cl"
PARENT_TARGET libspirv-builtins)

add_libclc_builtin_set(clc-${arch_suffix}
Expand All @@ -240,7 +247,7 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
FILES ${lib_files}
LIB_DEP libspirv-${arch_suffix}
ALIASES ${${d}_aliases}
GENERATE_TARGET "generate_convert_spirv.cl"
GENERATE_TARGET "generate_convert_clc.cl"
PARENT_TARGET libclc-builtins)
endforeach( d )
endforeach( t )
Expand Down
4 changes: 2 additions & 2 deletions libclc/cmake/modules/AddLibclc.cmake
View file
Original file line number Diff line number Diff line change
Expand Up @@ -42,8 +42,8 @@ endfunction(add_libclc_alias alias target)
macro(add_libclc_builtin_set arch_suffix)
cmake_parse_arguments(ARG
""
"TRIPLE;TARGET_ENV;LIB_DEP;GENERATE_TARGET;PARENT_TARGET"
"FILES;ALIASES;COMPILE_OPT"
"TRIPLE;TARGET_ENV;LIB_DEP;PARENT_TARGET"
"FILES;ALIASES;GENERATE_TARGET;COMPILE_OPT"
${ARGN})

if (DEFINED ${ARG_LIB_DEP})
Expand Down
Loading

0 comments on commit fc3b63d

Please sign in to comment.