Skip to content

Commit 52d3006

Browse files
authored
[SYCLomatic] Fix tool abort when __launch_bounds__ uses maxBlocksPerCluster (#2828)
Signed-off-by: Jiang, Zhiwei <zhiwei.jiang@intel.com>
1 parent e251b18 commit 52d3006

File tree

2 files changed

+11
-0
lines changed

2 files changed

+11
-0
lines changed

clang/lib/Sema/SemaDeclAttr.cpp

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -5464,10 +5464,16 @@ Sema::CreateLaunchBoundsAttr(const AttributeCommonInfo &CI, Expr *MaxThreads,
54645464

54655465
if (MaxBlocks) {
54665466
// '.maxclusterrank' ptx directive requires .target sm_90 or higher.
5467+
#ifdef SYCLomatic_CUSTOMIZATION
5468+
auto SM = OffloadArch::UNKNOWN;
5469+
#else
54675470
auto SM = SYCL().getOffloadArch(Context.getTargetInfo());
5471+
#endif // SYCLomatic_CUSTOMIZATION
54685472
if (SM == OffloadArch::UNKNOWN || SM < OffloadArch::SM_90) {
5473+
#ifndef SYCLomatic_CUSTOMIZATION
54695474
Diag(MaxBlocks->getBeginLoc(), diag::warn_cuda_maxclusterrank_sm_90)
54705475
<< OffloadArchToString(SM) << CI << MaxBlocks->getSourceRange();
5476+
#endif // SYCLomatic_CUSTOMIZATION
54715477
// Ignore it by setting MaxBlocks to null;
54725478
MaxBlocks = nullptr;
54735479
} else {

clang/test/dpct/launch_bounds.cu

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,5 @@
1+
// RUN: dpct --format-range=none --out-root %T/launch_bounds %s --cuda-include-path="%cuda-path/include"
2+
// RUN: FileCheck %s --match-full-lines --input-file %T/launch_bounds/launch_bounds.dp.cpp
3+
4+
// CHECK: void foo() {}
5+
__global__ void __launch_bounds__(32, 1, 1) foo() {}

0 commit comments

Comments
 (0)