Skip to content

Commit c16b5cf

Browse files
victor-edsvictor-eds
committed
Apply changes to SYCL types
1 parent a562e0b commit c16b5cf

File tree

1 file changed

+17
-17
lines changed

1 file changed

+17
-17
lines changed

‎polygeist/tools/cgeist/Lib/TypeUtils.cc

Lines changed: 17 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -176,7 +176,7 @@ mlir::Type getSYCLType(const clang::RecordType *RT,
176176
llvm::SmallVector<mlir::Type, 4> Body;
177177

178178
for (const auto *Field : RD->fields())
179-
Body.push_back(CGT.getMLIRType(Field->getType()));
179+
Body.push_back(CGT.getMLIRTypeForMem(Field->getType()));
180180

181181
if (const auto *CTS =
182182
llvm::dyn_cast<clang::ClassTemplateSpecializationDecl>(RD)) {
@@ -193,7 +193,7 @@ mlir::Type getSYCLType(const clang::RecordType *RT,
193193
}
194194
case TypeEnum::Accessor: {
195195
const auto Type =
196-
CGT.getMLIRType(CTS->getTemplateArgs().get(0).getAsType());
196+
CGT.getMLIRTypeForMem(CTS->getTemplateArgs().get(0).getAsType());
197197
const auto Dim =
198198
CTS->getTemplateArgs().get(1).getAsIntegral().getExtValue();
199199
const auto MemAccessMode = static_cast<mlir::sycl::MemoryAccessMode>(
@@ -206,7 +206,7 @@ mlir::Type getSYCLType(const clang::RecordType *RT,
206206
// TODO: we should push the non-empty base classes in a more general way.
207207
if (MemTargetMode == mlir::sycl::MemoryTargetMode::Local) {
208208
assert(Body.empty());
209-
Body.push_back(CGT.getMLIRType(CTS->bases_begin()->getType()));
209+
Body.push_back(CGT.getMLIRTypeForMem(CTS->bases_begin()->getType()));
210210
}
211211

212212
return mlir::sycl::AccessorType::get(CGT.getModule()->getContext(), Type,
@@ -227,7 +227,7 @@ mlir::Type getSYCLType(const clang::RecordType *RT,
227227
}
228228
case TypeEnum::Atomic: {
229229
const auto Type =
230-
CGT.getMLIRType(CTS->getTemplateArgs().get(0).getAsType());
230+
CGT.getMLIRTypeForMem(CTS->getTemplateArgs().get(0).getAsType());
231231
const int AddrSpace =
232232
CTS->getTemplateArgs().get(1).getAsIntegral().getExtValue();
233233
return mlir::sycl::AtomicType::get(
@@ -236,12 +236,12 @@ mlir::Type getSYCLType(const clang::RecordType *RT,
236236
}
237237
case TypeEnum::GetOp: {
238238
const auto Type =
239-
CGT.getMLIRType(CTS->getTemplateArgs().get(0).getAsType());
239+
CGT.getMLIRTypeForMem(CTS->getTemplateArgs().get(0).getAsType());
240240
return mlir::sycl::GetOpType::get(CGT.getModule()->getContext(), Type);
241241
}
242242
case TypeEnum::GetScalarOp: {
243243
const auto Type =
244-
CGT.getMLIRType(CTS->getTemplateArgs().get(0).getAsType());
244+
CGT.getMLIRTypeForMem(CTS->getTemplateArgs().get(0).getAsType());
245245
return mlir::sycl::GetScalarOpType::get(CGT.getModule()->getContext(),
246246
Type, Body);
247247
}
@@ -260,7 +260,7 @@ mlir::Type getSYCLType(const clang::RecordType *RT,
260260
case TypeEnum::ID: {
261261
const auto Dim =
262262
CTS->getTemplateArgs().get(0).getAsIntegral().getExtValue();
263-
Body.push_back(CGT.getMLIRType(CTS->bases_begin()->getType()));
263+
Body.push_back(CGT.getMLIRTypeForMem(CTS->bases_begin()->getType()));
264264
return mlir::sycl::IDType::get(CGT.getModule()->getContext(), Dim, Body);
265265
}
266266
case TypeEnum::ItemBase: {
@@ -287,7 +287,7 @@ mlir::Type getSYCLType(const clang::RecordType *RT,
287287
}
288288
case TypeEnum::LocalAccessorBase: {
289289
const auto Type =
290-
CGT.getMLIRType(CTS->getTemplateArgs().get(0).getAsType());
290+
CGT.getMLIRTypeForMem(CTS->getTemplateArgs().get(0).getAsType());
291291
const auto Dim =
292292
CTS->getTemplateArgs().get(1).getAsIntegral().getExtValue();
293293
const auto MemAccessMode = static_cast<mlir::sycl::MemoryAccessMode>(
@@ -297,10 +297,10 @@ mlir::Type getSYCLType(const clang::RecordType *RT,
297297
}
298298
case TypeEnum::LocalAccessor: {
299299
const auto Type =
300-
CGT.getMLIRType(CTS->getTemplateArgs().get(0).getAsType());
300+
CGT.getMLIRTypeForMem(CTS->getTemplateArgs().get(0).getAsType());
301301
const auto Dim =
302302
CTS->getTemplateArgs().get(1).getAsIntegral().getExtValue();
303-
Body.push_back(CGT.getMLIRType(CTS->bases_begin()->getType()));
303+
Body.push_back(CGT.getMLIRTypeForMem(CTS->bases_begin()->getType()));
304304
return mlir::sycl::LocalAccessorType::get(CGT.getModule()->getContext(),
305305
Type, Dim, Body);
306306
}
@@ -316,7 +316,7 @@ mlir::Type getSYCLType(const clang::RecordType *RT,
316316
}
317317
case TypeEnum::MultiPtr: {
318318
const auto Type =
319-
CGT.getMLIRType(CTS->getTemplateArgs().get(0).getAsType());
319+
CGT.getMLIRTypeForMem(CTS->getTemplateArgs().get(0).getAsType());
320320
const int AddrSpace =
321321
CTS->getTemplateArgs().get(1).getAsIntegral().getExtValue();
322322
const int DecAccess =
@@ -343,36 +343,36 @@ mlir::Type getSYCLType(const clang::RecordType *RT,
343343
case TypeEnum::Range: {
344344
const auto Dim =
345345
CTS->getTemplateArgs().get(0).getAsIntegral().getExtValue();
346-
Body.push_back(CGT.getMLIRType(CTS->bases_begin()->getType()));
346+
Body.push_back(CGT.getMLIRTypeForMem(CTS->bases_begin()->getType()));
347347
return mlir::sycl::RangeType::get(CGT.getModule()->getContext(), Dim,
348348
Body);
349349
}
350350
case TypeEnum::TupleCopyAssignableValueHolder: {
351351
const auto Type =
352-
CGT.getMLIRType(CTS->getTemplateArgs().get(0).getAsType());
352+
CGT.getMLIRTypeForMem(CTS->getTemplateArgs().get(0).getAsType());
353353
const auto IsTriviallyCopyAssignable =
354354
CTS->getTemplateArgs().get(1).getAsIntegral().getExtValue();
355-
Body.push_back(CGT.getMLIRType(CTS->bases_begin()->getType()));
355+
Body.push_back(CGT.getMLIRTypeForMem(CTS->bases_begin()->getType()));
356356
return mlir::sycl::TupleCopyAssignableValueHolderType::get(
357357
CGT.getModule()->getContext(), Type, IsTriviallyCopyAssignable, Body);
358358
}
359359
case TypeEnum::TupleValueHolder: {
360360
const auto Type =
361-
CGT.getMLIRType(CTS->getTemplateArgs().get(0).getAsType());
361+
CGT.getMLIRTypeForMem(CTS->getTemplateArgs().get(0).getAsType());
362362
return mlir::sycl::TupleValueHolderType::get(
363363
CGT.getModule()->getContext(), Type, Body);
364364
}
365365
case TypeEnum::Vec: {
366366
const auto ElemType =
367-
CGT.getMLIRType(CTS->getTemplateArgs().get(0).getAsType());
367+
CGT.getMLIRTypeForMem(CTS->getTemplateArgs().get(0).getAsType());
368368
const auto NumElems =
369369
CTS->getTemplateArgs().get(1).getAsIntegral().getExtValue();
370370
return mlir::sycl::VecType::get(CGT.getModule()->getContext(), ElemType,
371371
NumElems, Body);
372372
}
373373
case TypeEnum::SwizzleOp: {
374374
const auto VecType =
375-
CGT.getMLIRType(CTS->getTemplateArgs().get(0).getAsType())
375+
CGT.getMLIRTypeForMem(CTS->getTemplateArgs().get(0).getAsType())
376376
.cast<mlir::sycl::VecType>();
377377
const auto IndexesArgs = CTS->getTemplateArgs().get(4).getPackAsArray();
378378
SmallVector<int> Indexes;

0 commit comments

Comments
 (0)