Skip to content

[SYCL] Do not treat sycl::half kernel argument specially #5812

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 1 commit into from
Mar 15, 2022
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
62 changes: 0 additions & 62 deletions clang/lib/Sema/SemaSYCL.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -92,10 +92,6 @@ class Util {

static bool isSyclSpecialType(QualType Ty);

/// Checks whether given clang type is a full specialization of the SYCL
/// half class.
static bool isSyclHalfType(QualType Ty);

/// Checks whether given clang type is a full specialization of the SYCL
/// accessor_property_list class.
static bool isAccessorPropertyListType(QualType Ty);
Expand Down Expand Up @@ -1313,8 +1309,6 @@ class KernelObjVisitor {
QualType FieldTy, HandlerTys &... Handlers) {
if (Util::isSyclSpecialType(FieldTy))
KF_FOR_EACH(handleSyclSpecialType, Field, FieldTy);
else if (Util::isSyclHalfType(FieldTy))
KF_FOR_EACH(handleSyclHalfType, Field, FieldTy);
else if (Util::isSyclSpecConstantType(FieldTy))
KF_FOR_EACH(handleSyclSpecConstantType, Field, FieldTy);
else if (FieldTy->isStructureOrClassType()) {
Expand Down Expand Up @@ -1383,11 +1377,6 @@ class SyclKernelFieldHandlerBase {
return true;
}

virtual bool handleSyclHalfType(const CXXRecordDecl *,
const CXXBaseSpecifier &, QualType) {
return true;
}
virtual bool handleSyclHalfType(FieldDecl *, QualType) { return true; }
virtual bool handleStructType(FieldDecl *, QualType) { return true; }
virtual bool handleUnionType(FieldDecl *, QualType) { return true; }
virtual bool handleReferenceType(FieldDecl *, QualType) { return true; }
Expand Down Expand Up @@ -1815,16 +1804,6 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler {
return true;
}

bool handleSyclHalfType(const CXXRecordDecl *, const CXXBaseSpecifier &,
QualType) final {
CollectionStack.back() = true;
return true;
}
bool handleSyclHalfType(FieldDecl *, QualType) final {
CollectionStack.back() = true;
return true;
}

bool handlePointerType(FieldDecl *, QualType) final {
CollectionStack.back() = true;
return true;
Expand Down Expand Up @@ -2230,11 +2209,6 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
return handleScalarType(FD, FieldTy);
}

bool handleSyclHalfType(FieldDecl *FD, QualType FieldTy) final {
addParam(FD, FieldTy);
return true;
}

// Generate kernel argument to initialize specialization constants.
void handleSyclKernelHandlerType() {
ASTContext &Context = SemaRef.getASTContext();
Expand All @@ -2251,7 +2225,6 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
return ArrayRef<ParmVarDecl *>(std::begin(Params) + LastParamIndex,
std::end(Params));
}
using SyclKernelFieldHandler::handleSyclHalfType;
};

class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler {
Expand Down Expand Up @@ -2328,12 +2301,6 @@ class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler {
bool handleUnionType(FieldDecl *FD, QualType FieldTy) final {
return handleScalarType(FD, FieldTy);
}

bool handleSyclHalfType(FieldDecl *FD, QualType FieldTy) final {
addParam(FieldTy);
return true;
}
using SyclKernelFieldHandler::handleSyclHalfType;
};

std::string getKernelArgDesc(StringRef KernelArgDescription) {
Expand Down Expand Up @@ -2476,16 +2443,10 @@ class SyclOptReportCreator : public SyclKernelFieldHandler {
return handleScalarType(FD, FieldTy);
}

bool handleSyclHalfType(FieldDecl *FD, QualType FieldTy) final {
addParam(FD, FieldTy);
return true;
}

void handleSyclKernelHandlerType() {
addParam(DC.getParamVarDeclsForCurrentField()[0]->getType(),
"SYCL2020 specialization constant");
}
using SyclKernelFieldHandler::handleSyclHalfType;
};

static CXXMethodDecl *getOperatorParens(const CXXRecordDecl *Rec) {
Expand Down Expand Up @@ -2940,11 +2901,6 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
return handleSpecialType(FD, Ty);
}

bool handleSyclHalfType(FieldDecl *FD, QualType Ty) final {
addSimpleFieldInit(FD, Ty);
return true;
}

bool handlePointerType(FieldDecl *FD, QualType FieldTy) final {
Expr *PointerRef =
createPointerParamReferenceExpr(FieldTy, StructDepth != 0);
Expand Down Expand Up @@ -3102,8 +3058,6 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
removeFieldMemberExpr(FD, ArrayType);
return true;
}

using SyclKernelFieldHandler::handleSyclHalfType;
};

// Kernels are only the unnamed-lambda feature if the feature is enabled, AND
Expand Down Expand Up @@ -3263,11 +3217,6 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
return handleScalarType(FD, FieldTy);
}

bool handleSyclHalfType(FieldDecl *FD, QualType FieldTy) final {
addParam(FD, FieldTy, SYCLIntegrationHeader::kind_std_layout);
return true;
}

void handleSyclKernelHandlerType(QualType Ty) {
// The compiler generated kernel argument used to initialize SYCL 2020
// specialization constants, `specialization_constants_buffer`, should
Expand Down Expand Up @@ -3321,7 +3270,6 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
}

using SyclKernelFieldHandler::enterStruct;
using SyclKernelFieldHandler::handleSyclHalfType;
using SyclKernelFieldHandler::leaveStruct;
};

Expand Down Expand Up @@ -5158,16 +5106,6 @@ bool Util::isSyclSpecialType(const QualType Ty) {
return RecTy->hasAttr<SYCLSpecialClassAttr>();
}

bool Util::isSyclHalfType(QualType Ty) {
std::array<DeclContextDesc, 5> Scopes = {
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "detail"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "half_impl"),
Util::MakeDeclContextDesc(Decl::Kind::CXXRecord, "half")};
return matchQualifiedTypeName(Ty, Scopes);
}

bool Util::isSyclSpecConstantType(QualType Ty) {
std::array<DeclContextDesc, 6> Scopes = {
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
Expand Down
4 changes: 2 additions & 2 deletions clang/test/SemaSYCL/decomposition.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -144,12 +144,12 @@ int main() {
myQueue.submit([&](sycl::handler &h) {
h.single_task<class Half1>([=]() { return t1.i; });
});
// CHECK: FunctionDecl {{.*}}Half1{{.*}} 'void (sycl::half, sycl::half, sycl::half, StructNonDecomposed, int)'
// CHECK: FunctionDecl {{.*}}Half1{{.*}} 'void (StructWithArray<StructWithHalf>)'

DerivedStruct<StructWithHalf> t2;
myQueue.submit([&](sycl::handler &h) {
h.single_task<class Half2>([=]() { return t2.i; });
});
// CHECK: FunctionDecl {{.*}}Half2{{.*}} 'void (sycl::half, StructNonDecomposed, int)'
// CHECK: FunctionDecl {{.*}}Half2{{.*}} 'void (DerivedStruct<StructWithHalf>)'
}
}