Skip to content

Commit 1b8714a

Browse files
authored
[SYCLomatic] Fix migration of initializer of __constant__ array with incomplete size (#704)
Signed-off-by: Cai, Justin <[email protected]>
1 parent 079c40e commit 1b8714a

File tree

2 files changed

+16
-3
lines changed

2 files changed

+16
-3
lines changed

clang/lib/DPCT/AnalysisInfo.h

Lines changed: 11 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -2197,6 +2197,15 @@ class CtTypeInfo {
21972197
// will follow as comments. If NeedSizeFold is false, original size expression
21982198
// will be the size string.
21992199
CtTypeInfo(const TypeLoc &TL, bool NeedSizeFold = false);
2200+
CtTypeInfo(const VarDecl *D, bool NeedSizeFold = false)
2201+
: CtTypeInfo(D->getTypeSourceInfo()->getTypeLoc(), NeedSizeFold) {
2202+
if (D &&
2203+
D->getTypeSourceInfo()
2204+
->getTypeLoc().getTypeLocClass()
2205+
== TypeLoc::IncompleteArray)
2206+
if (auto CAT = dyn_cast<ConstantArrayType>(D->getType()))
2207+
Range[0] = std::to_string(CAT->getSize().getZExtValue());
2208+
}
22002209

22012210
inline const std::string &getBaseName() { return BaseName; }
22022211

@@ -2294,10 +2303,9 @@ class CtTypeInfo {
22942303
class VarInfo {
22952304
public:
22962305
VarInfo(unsigned Offset, const std::string &FilePathIn,
2297-
const DeclaratorDecl *Var, bool NeedFoldSize = false)
2306+
const VarDecl *Var, bool NeedFoldSize = false)
22982307
: FilePath(FilePathIn), Offset(Offset), Name(Var->getName()),
2299-
Ty(std::make_shared<CtTypeInfo>(Var->getTypeSourceInfo()->getTypeLoc(),
2300-
NeedFoldSize)) {}
2308+
Ty(std::make_shared<CtTypeInfo>(Var, NeedFoldSize)) {}
23012309

23022310
inline const std::string &getFilePath() { return FilePath; }
23032311
inline unsigned getOffset() { return Offset; }

clang/test/dpct/cuda_const.cu

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -31,8 +31,13 @@ __constant__ int *const_ptr;
3131

3232
// CHECK: static dpct::constant_memory<int, 1> const_init(sycl::range<1>(5), {1, 2, 3, 7, 8});
3333
__constant__ int const_init[5] = {1, 2, 3, 7, 8};
34+
// CHECK: static dpct::constant_memory<int, 1> incomplete_size_init(sycl::range<1>(5), {1, 2, 3, 7, 8});
35+
__constant__ int incomplete_size_init[] = {1, 2, 3, 7, 8};
3436
// CHECK: static dpct::constant_memory<int, 2> const_init_2d(sycl::range<2>(5, 5), {{[{][{]}}1, 2, 3, 7, 8}, {2, 4, 5, 8, 2}, {4, 7, 8, 0}, {1, 3}, {4, 0, 56}});
3537
__constant__ int const_init_2d[5][5] = {{1, 2, 3, 7, 8}, {2, 4, 5, 8, 2}, {4, 7, 8, 0}, {1, 3}, {4, 0, 56}};
38+
// CHECK: static dpct::constant_memory<int, 2> incomplete_size_init_2d(sycl::range<2>(3, 2), { {1,2},{3,4},{5,6}});
39+
__constant__ int incomplete_size_init_2d[][2] = { {1,2},{3,4},{5,6}};
40+
3641

3742
// CHECK: struct FuncObj {
3843
// CHECK-NEXT: void operator()(float *out, int index, float *const_angle) {

0 commit comments

Comments
 (0)