Skip to content
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.

Sign up for GitHub

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

Jump to bottom

[SYCL][CUDA] Enable sub-group loads and stores #2763

Merged
merged 1 commit into from
Jan 13, 2021
Merged

[SYCL][CUDA] Enable sub-group loads and stores #2763

merged 1 commit into from
Jan 13, 2021

Conversation

Pennycook
Copy link
Contributor

While implementing this I uncovered a bug in the load_store test, which assumed that it would always be safe to execute a sub-group load or store that accessed invalid memory. Admittedly this is something that the specification is unclear about, and something that we should revisit after some further discussion.

In my tests, everything works except for half. Trying to compile the test for half causes the compiler itself to crash, and I don't understand why, which is why I've opened this as a draft PR for now. @bader, can you please recommend somebody to take a look at this error and help fix it?

clang-12: /home/sjpennyc/llvm/clang/include/clang/AST/Type.h:387: void clang::Qualifiers::addAddressSpace(clang::LangAS): Assertion `space != LangAS::Default' failed.
PLEASE submit a bug report to https://bugs.llvm.org/ and include the crash backtrace, preprocessed source, and associated run script.
Stack dump:
0.      Program arguments: /home/sjpennyc/llvm/build/bin/clang-12 -cc1 -triple nvptx64-nvidia-cuda-sycldevice -fsycl -fsycl-is-device -fdeclare-spirv-builtins -mllvm -sycl-opt -aux-triple x86_64-unknown-linux-gnu -Wno-sycl-strict -sycl-std=2020 -emit-llvm-bc -emit-llvm-uselists -disable-free -main-file-name load_store.cpp -mrelocation-model static -mframe-pointer=all -fno-rounding-math -fno-verbose-asm -no-integrated-as -aux-target-cpu x86-64 -internal-isystem /home/sjpennyc/llvm/build//bin/../include/sycl -mlink-builtin-bitcode /home/sjpennyc/llvm/build/lib/clang/12.0.0/../../clc/libspirv-nvptx64--nvidiacl.bc -mlink-builtin-bitcode /usr/local/cuda/nvvm/libdevice/libdevice.10.bc -target-feature +ptx70 -target-sdk-version=11.0 -target-cpu sm_50 -fno-split-dwarf-inlining -debugger-tuning=gdb -resource-dir /home/sjpennyc/llvm/build/lib/clang/12.0.0 -internal-isystem /home/sjpennyc/llvm/build//bin/../include/sycl -I . -I/home/sjpennyc/llvm/build//lib/clang/11.0.0/include/ -I/home/sjpennyc/llvm/build//lib/clang/11.0.0/include/ -I. -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/x86_64-linux-gnu/c++/9 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/x86_64-linux-gnu/c++/9 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/backward -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/x86_64-linux-gnu/c++/9 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/x86_64-linux-gnu/c++/9 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/backward -internal-isystem /usr/local/include -internal-isystem /home/sjpennyc/llvm/build/lib/clang/12.0.0/include -internal-externc-isystem /usr/include/x86_64-linux-gnu -internal-externc-isystem /include -internal-externc-isystem /usr/include -internal-isystem /usr/local/include -internal-isystem /home/sjpennyc/llvm/build/lib/clang/12.0.0/include -internal-externc-isystem /usr/include/x86_64-linux-gnu -internal-externc-isystem /include -internal-externc-isystem /usr/include -std=c++17 -fdeprecated-macro -fno-dwarf-directory-asm -fdebug-compilation-dir /home/sjpennyc/llvm/sycl/test/on-device/sub_group -ferror-limit 19 -fgnuc-version=4.2.1 -fcxx-exceptions -fexceptions -fcolor-diagnostics -o /tmp/load_store-f29d19.bc -x c++ load_store.cpp
1.      <eof> parser at end of file
2.      load_store.cpp:25:35: instantiating function definition 'check<cl::sycl::detail::half_impl::half, 1>'
3.      /home/sjpennyc/llvm/build//bin/../include/sycl/CL/sycl/ONEAPI/sub_group.hpp:304:3: instantiating function definition 'cl::sycl::ONEAPI::sub_group::load<1, cl::sycl::detail::half_impl::half, cl::sycl::access::address_space::global_space>'
4.      /home/sjpennyc/llvm/build//bin/../include/sycl/CL/sycl/types.hpp:591:30: instantiating class definition 'cl::sycl::vec<cl::sycl::detail::half_impl::half, 1>::conjunction<cl::sycl::vec<cl::sycl::detail::half_impl::half, 1>::TypeChecker<__global cl::sycl::detail::half_impl::half, cl::sycl::detail::half_impl::half>>'
5.      /home/sjpennyc/llvm/build//bin/../include/sycl/CL/sycl/types.hpp:598:10: instantiating class definition 'cl::sycl::vec<cl::sycl::detail::half_impl::half, 1>::TypeChecker<__global cl::sycl::detail::half_impl::half, cl::sycl::detail::half_impl::half>'
6.      /usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/type_traits:1335:12: instantiating class definition 'std::is_convertible<__global cl::sycl::detail::half_impl::half, cl::sycl::detail::half_impl::half>'
7.      /usr/lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/type_traits:1308:12: instantiating class definition 'std::__is_convertible_helper<__global cl::sycl::detail::half_impl::half, cl::sycl::detail::half_impl::half, false>'
  #0 0x00005581a463e58c llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/home/sjpennyc/llvm/build/bin/clang-12+0x1fb258c)
  #1 0x00005581a463c264 llvm::sys::RunSignalHandlers() (/home/sjpennyc/llvm/build/bin/clang-12+0x1fb0264)
  #2 0x00005581a463c3d3 SignalHandler(int) (/home/sjpennyc/llvm/build/bin/clang-12+0x1fb03d3)
  #3 0x00007f30353f5980 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x12980)
  #4 0x00007f303405bfb7 raise /build/glibc-S7xCS9/glibc-2.27/signal/../sysdeps/unix/sysv/linux/raise.c:51:0
  #5 0x00007f303405d921 abort /build/glibc-S7xCS9/glibc-2.27/stdlib/abort.c:81:0
  #6 0x00007f303404d48a __assert_fail_base /build/glibc-S7xCS9/glibc-2.27/assert/assert.c:89:0
  #7 0x00007f303404d502 (/lib/x86_64-linux-gnu/libc.so.6+0x30502)
  #8 0x00005581a6668d27 TryReferenceInitializationCore(clang::Sema&, clang::InitializedEntity const&, clang::InitializationKind const&, clang::Expr*, clang::QualType, clang::QualType, clang::Qualifiers, clang::QualType, clang::QualType, clang::Qualifiers, clang::InitializationSequence&) (/home/sjpennyc/llvm/build/bin/clang-12+0x3fdcd27)
  #9 0x00005581a667311f clang::InitializationSequence::InitializeFrom(clang::Sema&, clang::InitializedEntity const&, clang::InitializationKind const&, llvm::MutableArrayRef<clang::Expr*>, bool, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x3fe711f)
 #10 0x00005581a667c844 clang::Sema::PerformCopyInitialization(clang::InitializedEntity const&, clang::SourceLocation, clang::ActionResult<clang::Expr*, true>, bool, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x3ff0844)
 #11 0x00005581a65294b0 clang::Sema::GatherArgumentsForCall(clang::SourceLocation, clang::FunctionDecl*, clang::FunctionProtoType const*, unsigned int, llvm::ArrayRef<clang::Expr*>, llvm::SmallVectorImpl<clang::Expr*>&, clang::Sema::VariadicCallType, bool, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x3e9d4b0)
 #12 0x00005581a6415240 clang::Sema::CompleteConstructorCall(clang::CXXConstructorDecl*, llvm::MutableArrayRef<clang::Expr*>, clang::SourceLocation, llvm::SmallVectorImpl<clang::Expr*>&, bool, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x3d89240)
 #13 0x00005581a6666069 PerformConstructorInitialization(clang::Sema&, clang::InitializedEntity const&, clang::InitializationKind const&, llvm::MutableArrayRef<clang::Expr*>, clang::InitializationSequence::Step const&, bool&, bool, bool, clang::SourceLocation, clang::SourceLocation) (/home/sjpennyc/llvm/build/bin/clang-12+0x3fda069)
 #14 0x00005581a6677a05 clang::InitializationSequence::Perform(clang::Sema&, clang::InitializedEntity const&, clang::InitializationKind const&, llvm::MutableArrayRef<clang::Expr*>, clang::QualType*) (/home/sjpennyc/llvm/build/bin/clang-12+0x3feba05)
 #15 0x00005581a667cb72 clang::Sema::PerformCopyInitialization(clang::InitializedEntity const&, clang::SourceLocation, clang::ActionResult<clang::Expr*, true>, bool, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x3ff0b72)
 #16 0x00005581a65294b0 clang::Sema::GatherArgumentsForCall(clang::SourceLocation, clang::FunctionDecl*, clang::FunctionProtoType const*, unsigned int, llvm::ArrayRef<clang::Expr*>, llvm::SmallVectorImpl<clang::Expr*>&, clang::Sema::VariadicCallType, bool, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x3e9d4b0)
 #17 0x00005581a6529f20 clang::Sema::ConvertArgumentsForCall(clang::CallExpr*, clang::Expr*, clang::FunctionDecl*, clang::FunctionProtoType const*, llvm::ArrayRef<clang::Expr*>, clang::SourceLocation, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x3e9df20)
 #18 0x00005581a652b75e clang::Sema::BuildResolvedCallExpr(clang::Expr*, clang::NamedDecl*, clang::SourceLocation, llvm::ArrayRef<clang::Expr*>, clang::SourceLocation, clang::Expr*, bool, clang::CallExpr::ADLCallKind) (/home/sjpennyc/llvm/build/bin/clang-12+0x3e9f75e)
 #19 0x00005581a67ac47d FinishOverloadedCallExpr(clang::Sema&, clang::Scope*, clang::Expr*, clang::UnresolvedLookupExpr*, clang::SourceLocation, llvm::MutableArrayRef<clang::Expr*>, clang::SourceLocation, clang::Expr*, clang::OverloadCandidateSet*, clang::OverloadCandidate**, clang::OverloadingResult, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x412047d)
 #20 0x00005581a67acee7 clang::Sema::BuildOverloadedCallExpr(clang::Scope*, clang::Expr*, clang::UnresolvedLookupExpr*, clang::SourceLocation, llvm::MutableArrayRef<clang::Expr*>, clang::SourceLocation, clang::Expr*, bool, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x4120ee7)
 #21 0x00005581a652d189 clang::Sema::BuildCallExpr(clang::Scope*, clang::Expr*, clang::SourceLocation, llvm::MutableArrayRef<clang::Expr*>, clang::SourceLocation, clang::Expr*, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x3ea1189)
 #22 0x00005581a652dd61 clang::Sema::ActOnCallExpr(clang::Scope*, clang::Expr*, clang::SourceLocation, llvm::MutableArrayRef<clang::Expr*>, clang::SourceLocation, clang::Expr*) (/home/sjpennyc/llvm/build/bin/clang-12+0x3ea1d61)
 #23 0x00005581a69b91c2 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformCallExpr(clang::CallExpr*) (/home/sjpennyc/llvm/build/bin/clang-12+0x432d1c2)
 #24 0x00005581a69b1874 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformExpr(clang::Expr*) (/home/sjpennyc/llvm/build/bin/clang-12+0x4325874)
 #25 0x00005581a69bbbc4 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformDecltypeType(clang::TypeLocBuilder&, clang::DecltypeTypeLoc) (/home/sjpennyc/llvm/build/bin/clang-12+0x432fbc4)
 #26 0x00005581a69bcedb clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformType(clang::TypeLocBuilder&, clang::TypeLoc) (/home/sjpennyc/llvm/build/bin/clang-12+0x4330edb)
 #27 0x00005581a69c18b2 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformType(clang::TypeSourceInfo*) (/home/sjpennyc/llvm/build/bin/clang-12+0x43358b2)
 #28 0x00005581a69c939e clang::Sema::SubstType(clang::TypeSourceInfo*, clang::MultiLevelTemplateArgumentList const&, clang::SourceLocation, clang::DeclarationName, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x433d39e)
 #29 0x00005581a68ade05 SubstDefaultTemplateArgument(clang::Sema&, clang::TemplateDecl*, clang::SourceLocation, clang::SourceLocation, clang::TemplateTypeParmDecl*, llvm::SmallVectorImpl<clang::TemplateArgument>&) (/home/sjpennyc/llvm/build/bin/clang-12+0x4221e05)
 #30 0x00005581a68af64e clang::Sema::SubstDefaultTemplateArgumentIfAvailable(clang::TemplateDecl*, clang::SourceLocation, clang::SourceLocation, clang::Decl*, llvm::SmallVectorImpl<clang::TemplateArgument>&, bool&) (/home/sjpennyc/llvm/build/bin/clang-12+0x422364e)
 #31 0x00005581a6982fdb clang::Sema::FinishTemplateArgumentDeduction(clang::FunctionTemplateDecl*, llvm::SmallVectorImpl<clang::DeducedTemplateArgument>&, unsigned int, clang::FunctionDecl*&, clang::sema::TemplateDeductionInfo&, llvm::SmallVectorImpl<clang::Sema::OriginalCallArg> const*, bool, llvm::function_ref<bool ()>) (/home/sjpennyc/llvm/build/bin/clang-12+0x42f6fdb)
 #32 0x00005581a698426e void llvm::function_ref<void ()>::callback_fn<clang::Sema::DeduceTemplateArguments(clang::FunctionTemplateDecl*, clang::TemplateArgumentListInfo*, llvm::ArrayRef<clang::Expr*>, clang::FunctionDecl*&, clang::sema::TemplateDeductionInfo&, bool, llvm::function_ref<bool (llvm::ArrayRef<clang::QualType>)>)::'lambda1'()>(long) (/home/sjpennyc/llvm/build/bin/clang-12+0x42f826e)
 #33 0x00005581a6263a3f clang::Sema::runWithSufficientStackSpace(clang::SourceLocation, llvm::function_ref<void ()>) (/home/sjpennyc/llvm/build/bin/clang-12+0x3bd7a3f)
 #34 0x00005581a6992883 clang::Sema::DeduceTemplateArguments(clang::FunctionTemplateDecl*, clang::TemplateArgumentListInfo*, llvm::ArrayRef<clang::Expr*>, clang::FunctionDecl*&, clang::sema::TemplateDeductionInfo&, bool, llvm::function_ref<bool (llvm::ArrayRef<clang::QualType>)>) (/home/sjpennyc/llvm/build/bin/clang-12+0x4306883)
 #35 0x00005581a679f51b clang::Sema::AddTemplateOverloadCandidate(clang::FunctionTemplateDecl*, clang::DeclAccessPair, clang::TemplateArgumentListInfo*, llvm::ArrayRef<clang::Expr*>, clang::OverloadCandidateSet&, bool, bool, bool, clang::CallExpr::ADLCallKind, clang::OverloadCandidateParamOrder) (/home/sjpennyc/llvm/build/bin/clang-12+0x411351b)
 #36 0x00005581a67a068d AddOverloadedCallCandidate(clang::Sema&, clang::DeclAccessPair, clang::TemplateArgumentListInfo*, llvm::ArrayRef<clang::Expr*>, clang::OverloadCandidateSet&, bool, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x411468d)
 #37 0x00005581a67a0920 clang::Sema::AddOverloadedCallCandidates(clang::UnresolvedLookupExpr*, llvm::ArrayRef<clang::Expr*>, clang::OverloadCandidateSet&, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x4114920)
 #38 0x00005581a67a0cbc clang::Sema::buildOverloadedCallSet(clang::Scope*, clang::Expr*, clang::UnresolvedLookupExpr*, llvm::MutableArrayRef<clang::Expr*>, clang::SourceLocation, clang::OverloadCandidateSet*, clang::ActionResult<clang::Expr*, true>*) (/home/sjpennyc/llvm/build/bin/clang-12+0x4114cbc)
 #39 0x00005581a67ace79 clang::Sema::BuildOverloadedCallExpr(clang::Scope*, clang::Expr*, clang::UnresolvedLookupExpr*, clang::SourceLocation, llvm::MutableArrayRef<clang::Expr*>, clang::SourceLocation, clang::Expr*, bool, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x4120e79)
 #40 0x00005581a652d189 clang::Sema::BuildCallExpr(clang::Scope*, clang::Expr*, clang::SourceLocation, llvm::MutableArrayRef<clang::Expr*>, clang::SourceLocation, clang::Expr*, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x3ea1189)
 #41 0x00005581a652dd61 clang::Sema::ActOnCallExpr(clang::Scope*, clang::Expr*, clang::SourceLocation, llvm::MutableArrayRef<clang::Expr*>, clang::SourceLocation, clang::Expr*) (/home/sjpennyc/llvm/build/bin/clang-12+0x3ea1d61)
 #42 0x00005581a69b91c2 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformCallExpr(clang::CallExpr*) (/home/sjpennyc/llvm/build/bin/clang-12+0x432d1c2)
 #43 0x00005581a69b1874 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformExpr(clang::Expr*) (/home/sjpennyc/llvm/build/bin/clang-12+0x4325874)
 #44 0x00005581a69bbbc4 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformDecltypeType(clang::TypeLocBuilder&, clang::DecltypeTypeLoc) (/home/sjpennyc/llvm/build/bin/clang-12+0x432fbc4)
 #45 0x00005581a69bcedb clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformType(clang::TypeLocBuilder&, clang::TypeLoc) (/home/sjpennyc/llvm/build/bin/clang-12+0x4330edb)
 #46 0x00005581a69c18b2 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformType(clang::TypeSourceInfo*) (/home/sjpennyc/llvm/build/bin/clang-12+0x43358b2)
 #47 0x00005581a69c939e clang::Sema::SubstType(clang::TypeSourceInfo*, clang::MultiLevelTemplateArgumentList const&, clang::SourceLocation, clang::DeclarationName, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x433d39e)
 #48 0x00005581a69fcb72 clang::TemplateDeclInstantiator::InstantiateTypedefNameDecl(clang::TypedefNameDecl*, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x4370b72)
 #49 0x00005581a69fd11f clang::TemplateDeclInstantiator::VisitTypedefDecl(clang::TypedefDecl*) (/home/sjpennyc/llvm/build/bin/clang-12+0x437111f)
 #50 0x00005581a69ca37e clang::Sema::InstantiateClass(clang::SourceLocation, clang::CXXRecordDecl*, clang::CXXRecordDecl*, clang::MultiLevelTemplateArgumentList const&, clang::TemplateSpecializationKind, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x433e37e)
 #51 0x00005581a69e2ea7 clang::Sema::InstantiateClassTemplateSpecialization(clang::SourceLocation, clang::ClassTemplateSpecializationDecl*, clang::TemplateSpecializationKind, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x4356ea7)
 #52 0x00005581a6a28edf void llvm::function_ref<void ()>::callback_fn<clang::Sema::RequireCompleteTypeImpl(clang::SourceLocation, clang::QualType, clang::Sema::CompleteTypeKind, clang::Sema::TypeDiagnoser*)::'lambda'()>(long) (/home/sjpennyc/llvm/build/bin/clang-12+0x439cedf)
 #53 0x00005581a6263a3f clang::Sema::runWithSufficientStackSpace(clang::SourceLocation, llvm::function_ref<void ()>) (/home/sjpennyc/llvm/build/bin/clang-12+0x3bd7a3f)
 #54 0x00005581a6a325bb clang::Sema::RequireCompleteTypeImpl(clang::SourceLocation, clang::QualType, clang::Sema::CompleteTypeKind, clang::Sema::TypeDiagnoser*) (/home/sjpennyc/llvm/build/bin/clang-12+0x43a65bb)
 #55 0x00005581a6a329b5 clang::Sema::RequireCompleteType(clang::SourceLocation, clang::QualType, clang::Sema::CompleteTypeKind, clang::Sema::TypeDiagnoser&) (/home/sjpennyc/llvm/build/bin/clang-12+0x43a69b5)
 #56 0x00005581a6294a82 clang::Sema::RequireCompleteDeclContext(clang::CXXScopeSpec&, clang::DeclContext*) (/home/sjpennyc/llvm/build/bin/clang-12+0x3c08a82)
 #57 0x00005581a68be555 clang::Sema::CheckTypenameType(clang::ElaboratedTypeKeyword, clang::SourceLocation, clang::NestedNameSpecifierLoc, clang::IdentifierInfo const&, clang::SourceLocation, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x4232555)
 #58 0x00005581a69ab218 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::RebuildDependentNameType(clang::ElaboratedTypeKeyword, clang::SourceLocation, clang::NestedNameSpecifierLoc, clang::IdentifierInfo const*, clang::SourceLocation, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x431f218)
 #59 0x00005581a69c8d8a clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformDependentNameType(clang::TypeLocBuilder&, clang::DependentNameTypeLoc, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x433cd8a)
 #60 0x00005581a69bd1ce clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformType(clang::TypeLocBuilder&, clang::TypeLoc) (/home/sjpennyc/llvm/build/bin/clang-12+0x43311ce)
 #61 0x00005581a69c18b2 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformType(clang::TypeSourceInfo*) (/home/sjpennyc/llvm/build/bin/clang-12+0x43358b2)
 #62 0x00005581a69c939e clang::Sema::SubstType(clang::TypeSourceInfo*, clang::MultiLevelTemplateArgumentList const&, clang::SourceLocation, clang::DeclarationName, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x433d39e)
 #63 0x00005581a69c9808 clang::Sema::SubstBaseSpecifiers(clang::CXXRecordDecl*, clang::CXXRecordDecl*, clang::MultiLevelTemplateArgumentList const&) (/home/sjpennyc/llvm/build/bin/clang-12+0x433d808)
 #64 0x00005581a69ca20a clang::Sema::InstantiateClass(clang::SourceLocation, clang::CXXRecordDecl*, clang::CXXRecordDecl*, clang::MultiLevelTemplateArgumentList const&, clang::TemplateSpecializationKind, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x433e20a)
 #65 0x00005581a69e2ea7 clang::Sema::InstantiateClassTemplateSpecialization(clang::SourceLocation, clang::ClassTemplateSpecializationDecl*, clang::TemplateSpecializationKind, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x4356ea7)
 #66 0x00005581a6a28edf void llvm::function_ref<void ()>::callback_fn<clang::Sema::RequireCompleteTypeImpl(clang::SourceLocation, clang::QualType, clang::Sema::CompleteTypeKind, clang::Sema::TypeDiagnoser*)::'lambda'()>(long) (/home/sjpennyc/llvm/build/bin/clang-12+0x439cedf)
 #67 0x00005581a6263a3f clang::Sema::runWithSufficientStackSpace(clang::SourceLocation, llvm::function_ref<void ()>) (/home/sjpennyc/llvm/build/bin/clang-12+0x3bd7a3f)
 #68 0x00005581a6a325bb clang::Sema::RequireCompleteTypeImpl(clang::SourceLocation, clang::QualType, clang::Sema::CompleteTypeKind, clang::Sema::TypeDiagnoser*) (/home/sjpennyc/llvm/build/bin/clang-12+0x43a65bb)
 #69 0x00005581a6a329b5 clang::Sema::RequireCompleteType(clang::SourceLocation, clang::QualType, clang::Sema::CompleteTypeKind, clang::Sema::TypeDiagnoser&) (/home/sjpennyc/llvm/build/bin/clang-12+0x43a69b5)
 #70 0x00005581a645271f clang::Sema::CheckBaseSpecifier(clang::CXXRecordDecl*, clang::SourceRange, bool, clang::AccessSpecifier, clang::TypeSourceInfo*, clang::SourceLocation) (/home/sjpennyc/llvm/build/bin/clang-12+0x3dc671f)
 #71 0x00005581a69c984a clang::Sema::SubstBaseSpecifiers(clang::CXXRecordDecl*, clang::CXXRecordDecl*, clang::MultiLevelTemplateArgumentList const&) (/home/sjpennyc/llvm/build/bin/clang-12+0x433d84a)
 #72 0x00005581a69ca20a clang::Sema::InstantiateClass(clang::SourceLocation, clang::CXXRecordDecl*, clang::CXXRecordDecl*, clang::MultiLevelTemplateArgumentList const&, clang::TemplateSpecializationKind, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x433e20a)
 #73 0x00005581a69e2ea7 clang::Sema::InstantiateClassTemplateSpecialization(clang::SourceLocation, clang::ClassTemplateSpecializationDecl*, clang::TemplateSpecializationKind, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x4356ea7)
 #74 0x00005581a6a28edf void llvm::function_ref<void ()>::callback_fn<clang::Sema::RequireCompleteTypeImpl(clang::SourceLocation, clang::QualType, clang::Sema::CompleteTypeKind, clang::Sema::TypeDiagnoser*)::'lambda'()>(long) (/home/sjpennyc/llvm/build/bin/clang-12+0x439cedf)
 #75 0x00005581a6263a3f clang::Sema::runWithSufficientStackSpace(clang::SourceLocation, llvm::function_ref<void ()>) (/home/sjpennyc/llvm/build/bin/clang-12+0x3bd7a3f)
 #76 0x00005581a6a325bb clang::Sema::RequireCompleteTypeImpl(clang::SourceLocation, clang::QualType, clang::Sema::CompleteTypeKind, clang::Sema::TypeDiagnoser*) (/home/sjpennyc/llvm/build/bin/clang-12+0x43a65bb)
 #77 0x00005581a6a329b5 clang::Sema::RequireCompleteType(clang::SourceLocation, clang::QualType, clang::Sema::CompleteTypeKind, clang::Sema::TypeDiagnoser&) (/home/sjpennyc/llvm/build/bin/clang-12+0x43a69b5)
 #78 0x00005581a6294a82 clang::Sema::RequireCompleteDeclContext(clang::CXXScopeSpec&, clang::DeclContext*) (/home/sjpennyc/llvm/build/bin/clang-12+0x3c08a82)
 #79 0x00005581a64ffebb clang::Sema::BuildQualifiedDeclarationNameExpr(clang::CXXScopeSpec&, clang::DeclarationNameInfo const&, bool, clang::Scope const*, clang::TypeSourceInfo**) (/home/sjpennyc/llvm/build/bin/clang-12+0x3e73ebb)
 #80 0x00005581a69ce7ba clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformDependentScopeDeclRefExpr(clang::DependentScopeDeclRefExpr*, bool, clang::TypeSourceInfo**) (/home/sjpennyc/llvm/build/bin/clang-12+0x43427ba)
 #81 0x00005581a69b1967 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformExpr(clang::Expr*) (/home/sjpennyc/llvm/build/bin/clang-12+0x4325967)
 #82 0x00005581a69b238c clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformInitializer(clang::Expr*, bool) (.part.3935) (/home/sjpennyc/llvm/build/bin/clang-12+0x432638c)
 #83 0x00005581a69b3538 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformExprs(clang::Expr* const*, unsigned int, bool, llvm::SmallVectorImpl<clang::Expr*>&, bool*) (/home/sjpennyc/llvm/build/bin/clang-12+0x4327538)
 #84 0x00005581a69cc64d clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformCXXUnresolvedConstructExpr(clang::CXXUnresolvedConstructExpr*) (/home/sjpennyc/llvm/build/bin/clang-12+0x434064d)
 #85 0x00005581a69b1db3 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformExpr(clang::Expr*) (/home/sjpennyc/llvm/build/bin/clang-12+0x4325db3)
 #86 0x00005581a69c5a13 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformTemplateArgument(clang::TemplateArgumentLoc const&, clang::TemplateArgumentLoc&, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x4339a13)
 #87 0x00005581a69c7382 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformTemplateSpecializationType(clang::TypeLocBuilder&, clang::TemplateSpecializationTypeLoc, clang::TemplateName) (/home/sjpennyc/llvm/build/bin/clang-12+0x433b382)
 #88 0x00005581a69bd8eb clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformType(clang::TypeLocBuilder&, clang::TypeLoc) (/home/sjpennyc/llvm/build/bin/clang-12+0x43318eb)
 #89 0x00005581a69cfb8d clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformElaboratedType(clang::TypeLocBuilder&, clang::ElaboratedTypeLoc) (/home/sjpennyc/llvm/build/bin/clang-12+0x4343b8d)
 #90 0x00005581a69bcd1b clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformType(clang::TypeLocBuilder&, clang::TypeLoc) (/home/sjpennyc/llvm/build/bin/clang-12+0x4330d1b)
 #91 0x00005581a69c18b2 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformType(clang::TypeSourceInfo*) (/home/sjpennyc/llvm/build/bin/clang-12+0x43358b2)
 #92 0x00005581a69c939e clang::Sema::SubstType(clang::TypeSourceInfo*, clang::MultiLevelTemplateArgumentList const&, clang::SourceLocation, clang::DeclarationName, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x433d39e)
 #93 0x00005581a69c9808 clang::Sema::SubstBaseSpecifiers(clang::CXXRecordDecl*, clang::CXXRecordDecl*, clang::MultiLevelTemplateArgumentList const&) (/home/sjpennyc/llvm/build/bin/clang-12+0x433d808)
 #94 0x00005581a69ca20a clang::Sema::InstantiateClass(clang::SourceLocation, clang::CXXRecordDecl*, clang::CXXRecordDecl*, clang::MultiLevelTemplateArgumentList const&, clang::TemplateSpecializationKind, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x433e20a)
 #95 0x00005581a69e2ea7 clang::Sema::InstantiateClassTemplateSpecialization(clang::SourceLocation, clang::ClassTemplateSpecializationDecl*, clang::TemplateSpecializationKind, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x4356ea7)
 #96 0x00005581a6a28edf void llvm::function_ref<void ()>::callback_fn<clang::Sema::RequireCompleteTypeImpl(clang::SourceLocation, clang::QualType, clang::Sema::CompleteTypeKind, clang::Sema::TypeDiagnoser*)::'lambda'()>(long) (/home/sjpennyc/llvm/build/bin/clang-12+0x439cedf)
 #97 0x00005581a6263a3f clang::Sema::runWithSufficientStackSpace(clang::SourceLocation, llvm::function_ref<void ()>) (/home/sjpennyc/llvm/build/bin/clang-12+0x3bd7a3f)
 #98 0x00005581a6a325bb clang::Sema::RequireCompleteTypeImpl(clang::SourceLocation, clang::QualType, clang::Sema::CompleteTypeKind, clang::Sema::TypeDiagnoser*) (/home/sjpennyc/llvm/build/bin/clang-12+0x43a65bb)
 #99 0x00005581a6a329b5 clang::Sema::RequireCompleteType(clang::SourceLocation, clang::QualType, clang::Sema::CompleteTypeKind, clang::Sema::TypeDiagnoser&) (/home/sjpennyc/llvm/build/bin/clang-12+0x43a69b5)
#100 0x00005581a6294a82 clang::Sema::RequireCompleteDeclContext(clang::CXXScopeSpec&, clang::DeclContext*) (/home/sjpennyc/llvm/build/bin/clang-12+0x3c08a82)
#101 0x00005581a64ffebb clang::Sema::BuildQualifiedDeclarationNameExpr(clang::CXXScopeSpec&, clang::DeclarationNameInfo const&, bool, clang::Scope const*, clang::TypeSourceInfo**) (/home/sjpennyc/llvm/build/bin/clang-12+0x3e73ebb)
#102 0x00005581a69ce7ba clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformDependentScopeDeclRefExpr(clang::DependentScopeDeclRefExpr*, bool, clang::TypeSourceInfo**) (/home/sjpennyc/llvm/build/bin/clang-12+0x43427ba)
#103 0x00005581a69b1967 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformExpr(clang::Expr*) (/home/sjpennyc/llvm/build/bin/clang-12+0x4325967)
#104 0x00005581a69b1900 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformExpr(clang::Expr*) (/home/sjpennyc/llvm/build/bin/clang-12+0x4325900)
#105 0x00005581a69c5a13 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformTemplateArgument(clang::TemplateArgumentLoc const&, clang::TemplateArgumentLoc&, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x4339a13)
#106 0x00005581a69c7382 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformTemplateSpecializationType(clang::TypeLocBuilder&, clang::TemplateSpecializationTypeLoc, clang::TemplateName) (/home/sjpennyc/llvm/build/bin/clang-12+0x433b382)
#107 0x00005581a69c4ae5 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformTSIInObjectScope(clang::TypeLoc, clang::QualType, clang::NamedDecl*, clang::CXXScopeSpec&) (/home/sjpennyc/llvm/build/bin/clang-12+0x4338ae5)
#108 0x00005581a69c5631 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformNestedNameSpecifierLoc(clang::NestedNameSpecifierLoc, clang::QualType, clang::NamedDecl*) (/home/sjpennyc/llvm/build/bin/clang-12+0x4339631)
#109 0x00005581a69c8d54 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformDependentNameType(clang::TypeLocBuilder&, clang::DependentNameTypeLoc, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x433cd54)
#110 0x00005581a69bd1ce clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformType(clang::TypeLocBuilder&, clang::TypeLoc) (/home/sjpennyc/llvm/build/bin/clang-12+0x43311ce)
#111 0x00005581a69c18b2 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformType(clang::TypeSourceInfo*) (/home/sjpennyc/llvm/build/bin/clang-12+0x43358b2)
#112 0x00005581a69c194e clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformType(clang::QualType) (/home/sjpennyc/llvm/build/bin/clang-12+0x433594e)
#113 0x00005581a69c3c99 clang::Sema::SubstType(clang::QualType, clang::MultiLevelTemplateArgumentList const&, clang::SourceLocation, clang::DeclarationName) (/home/sjpennyc/llvm/build/bin/clang-12+0x4337c99)
#114 0x00005581a68c99d5 clang::Sema::CheckTemplateIdType(clang::TemplateName, clang::SourceLocation, clang::TemplateArgumentListInfo&) (/home/sjpennyc/llvm/build/bin/clang-12+0x423d9d5)
#115 0x00005581a69c7480 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformTemplateSpecializationType(clang::TypeLocBuilder&, clang::TemplateSpecializationTypeLoc, clang::TemplateName) (/home/sjpennyc/llvm/build/bin/clang-12+0x433b480)
#116 0x00005581a69bd8eb clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformType(clang::TypeLocBuilder&, clang::TypeLoc) (/home/sjpennyc/llvm/build/bin/clang-12+0x43318eb)
#117 0x00005581a69c18b2 clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformType(clang::TypeSourceInfo*) (/home/sjpennyc/llvm/build/bin/clang-12+0x43358b2)
#118 0x00005581a69c939e clang::Sema::SubstType(clang::TypeSourceInfo*, clang::MultiLevelTemplateArgumentList const&, clang::SourceLocation, clang::DeclarationName, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x433d39e)
#119 0x00005581a68ade05 SubstDefaultTemplateArgument(clang::Sema&, clang::TemplateDecl*, clang::SourceLocation, clang::SourceLocation, clang::TemplateTypeParmDecl*, llvm::SmallVectorImpl<clang::TemplateArgument>&) (/home/sjpennyc/llvm/build/bin/clang-12+0x4221e05)
#120 0x00005581a68af64e clang::Sema::SubstDefaultTemplateArgumentIfAvailable(clang::TemplateDecl*, clang::SourceLocation, clang::SourceLocation, clang::Decl*, llvm::SmallVectorImpl<clang::TemplateArgument>&, bool&) (/home/sjpennyc/llvm/build/bin/clang-12+0x422364e)
#121 0x00005581a6982fdb clang::Sema::FinishTemplateArgumentDeduction(clang::FunctionTemplateDecl*, llvm::SmallVectorImpl<clang::DeducedTemplateArgument>&, unsigned int, clang::FunctionDecl*&, clang::sema::TemplateDeductionInfo&, llvm::SmallVectorImpl<clang::Sema::OriginalCallArg> const*, bool, llvm::function_ref<bool ()>) (/home/sjpennyc/llvm/build/bin/clang-12+0x42f6fdb)
#122 0x00005581a698426e void llvm::function_ref<void ()>::callback_fn<clang::Sema::DeduceTemplateArguments(clang::FunctionTemplateDecl*, clang::TemplateArgumentListInfo*, llvm::ArrayRef<clang::Expr*>, clang::FunctionDecl*&, clang::sema::TemplateDeductionInfo&, bool, llvm::function_ref<bool (llvm::ArrayRef<clang::QualType>)>)::'lambda1'()>(long) (/home/sjpennyc/llvm/build/bin/clang-12+0x42f826e)
#123 0x00005581a6263a3f clang::Sema::runWithSufficientStackSpace(clang::SourceLocation, llvm::function_ref<void ()>) (/home/sjpennyc/llvm/build/bin/clang-12+0x3bd7a3f)
#124 0x00005581a6992883 clang::Sema::DeduceTemplateArguments(clang::FunctionTemplateDecl*, clang::TemplateArgumentListInfo*, llvm::ArrayRef<clang::Expr*>, clang::FunctionDecl*&, clang::sema::TemplateDeductionInfo&, bool, llvm::function_ref<bool (llvm::ArrayRef<clang::QualType>)>) (/home/sjpennyc/llvm/build/bin/clang-12+0x4306883)
#125 0x00005581a679f51b clang::Sema::AddTemplateOverloadCandidate(clang::FunctionTemplateDecl*, clang::DeclAccessPair, clang::TemplateArgumentListInfo*, llvm::ArrayRef<clang::Expr*>, clang::OverloadCandidateSet&, bool, bool, bool, clang::CallExpr::ADLCallKind, clang::OverloadCandidateParamOrder) (/home/sjpennyc/llvm/build/bin/clang-12+0x411351b)
#126 0x00005581a32abece TryUserDefinedConversion(clang::Sema&, clang::QualType, clang::InitializationKind const&, clang::Expr*, clang::InitializationSequence&, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0xc1fece)
#127 0x00005581a66733c8 clang::InitializationSequence::InitializeFrom(clang::Sema&, clang::InitializedEntity const&, clang::InitializationKind const&, llvm::MutableArrayRef<clang::Expr*>, bool, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x3fe73c8)
#128 0x00005581a667c844 clang::Sema::PerformCopyInitialization(clang::InitializedEntity const&, clang::SourceLocation, clang::ActionResult<clang::Expr*, true>, bool, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x3ff0844)
#129 0x00005581a67d5d95 clang::Sema::PerformMoveOrCopyInitialization(clang::InitializedEntity const&, clang::VarDecl const*, clang::QualType, clang::Expr*, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x4149d95)
#130 0x00005581a67ebeb2 clang::Sema::BuildReturnStmt(clang::SourceLocation, clang::Expr*) (/home/sjpennyc/llvm/build/bin/clang-12+0x415feb2)
#131 0x00005581a69e06bb clang::TreeTransform<(anonymous namespace)::TemplateInstantiator>::TransformCompoundStmt(clang::CompoundStmt*, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x43546bb)
#132 0x00005581a69e2c93 clang::Sema::SubstStmt(clang::Stmt*, clang::MultiLevelTemplateArgumentList const&) (/home/sjpennyc/llvm/build/bin/clang-12+0x4356c93)
#133 0x00005581a69f9f04 clang::Sema::InstantiateFunctionDefinition(clang::SourceLocation, clang::FunctionDecl*, bool, bool, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x436df04)
#134 0x00005581a69f8e5f clang::Sema::PerformPendingInstantiations(bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x436ce5f)
#135 0x00005581a69f9e87 clang::Sema::InstantiateFunctionDefinition(clang::SourceLocation, clang::FunctionDecl*, bool, bool, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x436de87)
#136 0x00005581a69f8e5f clang::Sema::PerformPendingInstantiations(bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x436ce5f)
#137 0x00005581a627a8ab clang::Sema::ActOnEndOfTranslationUnitFragment(clang::Sema::TUFragmentKind) (.part.1515) (/home/sjpennyc/llvm/build/bin/clang-12+0x3bee8ab)
#138 0x00005581a627afdf clang::Sema::ActOnEndOfTranslationUnit() (/home/sjpennyc/llvm/build/bin/clang-12+0x3beefdf)
#139 0x00005581a6178d96 clang::Parser::ParseTopLevelDecl(clang::OpaquePtr<clang::DeclGroupRef>&, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x3aecd96)
#140 0x00005581a616e4a9 clang::ParseAST(clang::Sema&, bool, bool) (/home/sjpennyc/llvm/build/bin/clang-12+0x3ae24a9)
#141 0x00005581a5517b70 clang::CodeGenAction::ExecuteAction() (/home/sjpennyc/llvm/build/bin/clang-12+0x2e8bb70)
#142 0x00005581a4ee6f89 clang::FrontendAction::Execute() (/home/sjpennyc/llvm/build/bin/clang-12+0x285af89)
#143 0x00005581a4ea097a clang::CompilerInstance::ExecuteAction(clang::FrontendAction&) (/home/sjpennyc/llvm/build/bin/clang-12+0x281497a)
#144 0x00005581a4fb2a76 clang::ExecuteCompilerInvocation(clang::CompilerInstance*) (/home/sjpennyc/llvm/build/bin/clang-12+0x2926a76)
#145 0x00005581a338827c cc1_main(llvm::ArrayRef<char const*>, char const*, void*) (/home/sjpennyc/llvm/build/bin/clang-12+0xcfc27c)
#146 0x00005581a3382cd9 ExecuteCC1Tool(llvm::SmallVectorImpl<char const*>&) (/home/sjpennyc/llvm/build/bin/clang-12+0xcf6cd9)
#147 0x00005581a32fc434 main (/home/sjpennyc/llvm/build/bin/clang-12+0xc70434)
#148 0x00007f303403ebf7 __libc_start_main /build/glibc-S7xCS9/glibc-2.27/csu/../csu/libc-start.c:344:0
#149 0x00005581a338283a _start (/home/sjpennyc/llvm/build/bin/clang-12+0xcf683a)

@Pennycook Pennycook added bug Something isn't working spec extension All issues/PRs related to extensions specifications cuda CUDA back-end labels Nov 11, 2020
@Pennycook Pennycook linked an issue Nov 11, 2020 that may be closed by this pull request
4 sub-group tests fail on CUDA #2754
Closed
@bader
Copy link
Contributor

bader commented Nov 11, 2020

can you please recommend somebody to take a look at this error and help fix it?

@premanandrao, could you help with that, please? The crash seems to happen in clang's Sema.

@premanandrao
Copy link
Contributor

can you please recommend somebody to take a look at this error and help fix it?

@premanandrao, could you help with that, please? The crash seems to happen in clang's Sema.

I will take a look.

@premanandrao premanandrao mentioned this pull request Nov 19, 2020
Merged
bader pushed a commit that referenced this pull request Dec 4, 2020
@premanandrao
[SYCL] Fix address space assertion with templates (#2798)
8905a8c 
In certain complex template instantiations, address space mismatches
cause the compiler to assert instead of issuing errors gracefully.
This fixes one such case identified in PR #2763.

Signed-off-by: Premanand M Rao <premanand.m.rao@intel.com>
@Pennycook Pennycook marked this pull request as ready for review December 17, 2020 21:11
@Pennycook Pennycook requested a review from a team as a code owner December 17, 2020 21:11
@Pennycook
Copy link
Contributor Author

#2798 appears to have fixed the tests, at least locally. Opening this up for review while GitHub repeats the tests.

@vladimirlaz
Copy link
Contributor

@Pennycook, could you please resolve conflicts. Please note that sub_group-related tests are in llvm-test-suite now which is available under Jenkins/Precommit.

@Pennycook
Copy link
Contributor Author

@vladimirlaz So I should move the test changes to llvm-test-suite? Should the changes to the tests be merged before or after the changes to the functionality?

@Pennycook
[SYCL][CUDA] Enable sub-group loads and stores
88f0935 
Rather than implementing the SubgroupBlockReadINTEL and SubgroupBlockWriteINTEL
intrinsics in libspirv, this commit adds a fallback path directly to the
sub-group header.  There are several reasons for this:

1) There are currently no INTEL extensions implemented in libspirv.
2) The load/store functions are expected to be rewritten soon to expose
   additional functionality, which may map to different SPIR-V.

Signed-off-by: John Pennycook <john.pennycook@intel.com>
@vladimirlaz
Copy link
Contributor

@vladimirlaz So I should move the test changes to llvm-test-suite? Should the changes to the tests be merged before or after the changes to the functionality?

PRs should be tested together. Submission should be done tests first and then product one after another.

@Pennycook Pennycook mentioned this pull request Jan 11, 2021
Merged
@vladimirlaz
Copy link
Contributor

The tests with product change produced expected results:
http://llvm-ci.intel.com:8080/job/SYCL_CI/job/intel/job/Lin/job/LLVM_Test_Suite/1269/display/redirect
http://llvm-ci.intel.com:8080/job/SYCL_CI/job/intel/job/Lin/job/LLVM_Test_Suite_CUDA/1014/display/redirect
http://llvm-ci.intel.com:8080/job/SYCL_CI/job/intel/job/Win/job/Test_Suite/462/display/redirect

@vladimirlaz vladimirlaz merged commit cfce996 into intel:sycl Jan 13, 2021
@Pennycook Pennycook deleted the cuda-sub-group-load-store branch January 28, 2021 18:23
jsji pushed a commit that referenced this pull request Nov 7, 2024
@delecui @jsji
Fix SPIR-V Builtin OpGenericCastToPtr transformation to OCL (#2763)
8659507 
OpGenericCastToPtrExplict is for dynamic cast and OpGenericCastToPtr is for static cast.

OpGenericCastToPtrExplict is already transformed to to_{global|local|private} OCL builtins, but the handling for OpGenericCastToPtr is missing.

Looks we can transform OpGenericCastToPtr to addrspacecast instruction directly in SPIRV to OCL transformation.

Signed-off-by: Cui, Dele <dele.cui@intel.com>

Original commit:
KhronosGroup/SPIRV-LLVM-Translator@21e96548e242860
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers

@vladimirlaz vladimirlaz vladimirlaz approved these changes

@rbegam rbegam Awaiting requested review from rbegam

Assignees
No one assigned
Labels
bug Something isn't working cuda CUDA back-end spec extension All issues/PRs related to extensions specifications
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.

4 sub-group tests fail on CUDA
4 participants
@Pennycook @bader @premanandrao @vladimirlaz