Add type converter for sycl::ArrayType (#46)

The runtime class of  `sycl::ArrayType`:
```
template <int dimensions = 1> class array {
…
  size_t common_array[dimensions];
…
}
```

Signed-off-by: Tsang, Whitney <whitney.tsang@intel.com>

Author: whitneywhtsang

mlir-sycl/lib/Conversion/SYCLToLLVM/SYCLToLLVM.cpp

@@ -50,6 +50,20 @@ static Optional<Type> getArrayTy(MLIRContext &context, unsigned dimNum,
 // Type conversion
 //===----------------------------------------------------------------------===//
 
+/// Converts SYCL array type to LLVM type.
+static Optional<Type> convertArrayType(sycl::ArrayType type,
+                                       LLVMTypeConverter &converter) {
+  assert(type.getBody().size() == 1 &&
+         "Expecting SYCL array body to have size 1");
+  assert(type.getBody()[0].isa<MemRefType>() &&
+         "Expecting SYCL array body entry to be MemRefType");
+  assert(type.getBody()[0].cast<MemRefType>().getElementType() ==
+             converter.getIndexType() &&
+         "Expecting SYCL array body entry element type to be the index type");
+  return getArrayTy(converter.getContext(), type.getDimension(),
+                    converter.getIndexType());
+}
+
 /// Converts SYCL range or id type to LLVM type, given \p dimNum - number of
 /// dimensions, \p name - the expected LLVM type name, \p converter - LLVM type
 /// converter.
@@ -127,8 +141,7 @@ void mlir::sycl::populateSYCLToLLVMTypeConversion(
     return convertAccessorType(type, typeConverter);
   });
   typeConverter.addConversion([&](sycl::ArrayType type) {
-    llvm_unreachable("SYCLToLLVM - sycl::ArrayType not handle (yet)");
-    return llvm::None;
+    return convertArrayType(type, typeConverter);
   });
   typeConverter.addConversion([&](sycl::GroupType type) {
     llvm_unreachable("SYCLToLLVM - sycl::GroupType not handle (yet)");

polygeist/tools/cgeist/Test/Verification/sycl/SYCLToLLVM/types.mlir

@@ -20,6 +20,8 @@
 // CHECK-DAG: [[ACCESSORIMPLDEVICE_2:%"class.cl::sycl::detail::AccessorImplDevice.*]] = type { [[ID_2]], [[RANGE_2]], [[RANGE_2]] }
 // CHECK-DAG: [[ACCESSOR_1:%"class.cl::sycl::accessor.*]] = type { [[ACCESSORIMPLDEVICE_1]], { i32 addrspace(1)* } }
 // CHECK-DAG: [[ACCESSOR_2:%"class.cl::sycl::accessor.*]] = type { [[ACCESSORIMPLDEVICE_2]], { i64 addrspace(1)* } }
+// CHECK: define void @test_array.1([[ARRAY_1]] %0)
+// CHECK: define void @test_array.2([[ARRAY_2]] %0)
 // CHECK: define void @test_id([[ID_1]] %0, [[ID_1]] %1)
 // CHECK: define void @test_range.1([[RANGE_1]] %0)
 // CHECK: define void @test_range.2([[RANGE_2]] %0)
@@ -28,6 +30,12 @@
 // CEHCK: define void @test_accessor.2([[ACCESSOR_2]] %0)
 
 module {
+  func.func @test_array.1(%arg0: !sycl.array<[1], (memref<1xi64>)>) {
+    return
+  }
+  func.func @test_array.2(%arg0: !sycl.array<[2], (memref<2xi64>)>) {
+    return
+  }
   func.func @test_id(%arg0: !sycl.id<1>, %arg1: !sycl.id<1>) {
     return
   }