[SYCL-MLIR][Opaque] Fix nd_item.get_group_linear_id lowering (#11016)

Use correct base type for `llvm.getelementptr` operation: it should be `nd_item`, not `group`.

---------

Signed-off-by: Victor Perez <[email protected]>

Author: victor-eds

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

@@ -1974,14 +1974,14 @@ class NDItemGetGroupLinearIDPattern
     // body is dropped and we can create group type more easily.
     const auto indices = GetMemberPattern<NDItemGroup>::getIndices();
     assert(indices.size() == 1 && "Expecting a single index");
-    const auto groupTy =
-        cast<NdItemType>(op.getNDItem().getType().getElementType())
-            .getBody()[indices[0]];
+    auto ndItemTy = cast<NdItemType>(op.getNDItem().getType().getElementType());
+    const auto groupTy = ndItemTy.getBody()[indices[0]];
     const auto convGroupTy = getTypeConverter()->convertType(groupTy);
+    Type convNDItemTy = getTypeConverter()->convertType(ndItemTy);
     bool useOpaquePointers = getTypeConverter()->useOpaquePointers();
     auto group = GetMemberPattern<NDItemGroup>::getRef(
-        rewriter, loc, convGroupTy, adaptor.getNDItem(), std::nullopt,
-        useOpaquePointers);
+        rewriter, loc, (useOpaquePointers) ? convNDItemTy : convGroupTy,
+        adaptor.getNDItem(), std::nullopt, useOpaquePointers);
     const auto thisTy = MemRefType::get(ShapedType::kDynamic, groupTy);
     // We have the already converted group, but, in order to not replicate
     // `sycl.group.get_group_linear_id` conversion to LLVM, we just reuse that

mlir-sycl/test/Conversion/SYCLToLLVM/sycl-methods-to-llvm.mlir

@@ -1033,7 +1033,7 @@ func.func @test_3(%nd: memref<?x!sycl_nd_item_3_>) -> i64 {
 
 // CHECK-LABEL:   llvm.func @test_1(
 // CHECK-SAME:                      %[[VAL_0:.*]]: !llvm.ptr) -> i64 {
-// CHECK:           %[[VAL_1:.*]] = llvm.getelementptr inbounds %[[VAL_0]][0, 2] : (!llvm.ptr) -> !llvm.ptr, !llvm.struct<"class.sycl::_V1::group.1", {{.*}}>
+// CHECK:           %[[VAL_1:.*]] = llvm.getelementptr inbounds %[[VAL_0]][0, 2] : (!llvm.ptr) -> !llvm.ptr, !llvm.struct<"class.sycl::_V1::nd_item.1", {{.*}}>
 // CHECK:           %[[VAL_3:.*]] = llvm.getelementptr inbounds %[[VAL_1]][0, 3, 0, 0, 0] : (!llvm.ptr) -> !llvm.ptr, !llvm.struct<"class.sycl::_V1::group.1", {{.*}}>
 // CHECK:           %[[VAL_4:.*]] = llvm.load %[[VAL_3]] : !llvm.ptr -> i64
 // CHECK:           llvm.return %[[VAL_4]] : i64
@@ -1045,7 +1045,7 @@ func.func @test_1(%nd: memref<?x!sycl_nd_item_1_>) -> i64 {
 
 // CHECK-LABEL:   llvm.func @test_2(
 // CHECK-SAME:                      %[[VAL_0:.*]]: !llvm.ptr) -> i64 {
-// CHECK:           %[[VAL_1:.*]] = llvm.getelementptr inbounds %[[VAL_0]][0, 2] : (!llvm.ptr) -> !llvm.ptr, !llvm.struct<"class.sycl::_V1::group.2", {{.*}}>
+// CHECK:           %[[VAL_1:.*]] = llvm.getelementptr inbounds %[[VAL_0]][0, 2] : (!llvm.ptr) -> !llvm.ptr, !llvm.struct<"class.sycl::_V1::nd_item.2", {{.*}}>
 // CHECK:           %[[VAL_3:.*]] = llvm.getelementptr inbounds %[[VAL_1]][0, 3, 0, 0, 0] : (!llvm.ptr) -> !llvm.ptr, !llvm.struct<"class.sycl::_V1::group.2", {{.*}}>
 // CHECK:           %[[VAL_4:.*]] = llvm.load %[[VAL_3]] : !llvm.ptr -> i64
 // CHECK:           %[[VAL_5:.*]] = llvm.getelementptr inbounds %[[VAL_1]][0, 2, 0, 0, 1] : (!llvm.ptr) -> !llvm.ptr, !llvm.struct<"class.sycl::_V1::group.2", {{.*}}>
@@ -1063,7 +1063,7 @@ func.func @test_2(%nd: memref<?x!sycl_nd_item_2_>) -> i64 {
 
 // CHECK-LABEL:   llvm.func @test_3(
 // CHECK-SAME:                      %[[VAL_0:.*]]: !llvm.ptr) -> i64 {
-// CHECK:           %[[VAL_1:.*]] = llvm.getelementptr inbounds %[[VAL_0]][0, 2] : (!llvm.ptr) -> !llvm.ptr, !llvm.struct<"class.sycl::_V1::group.3", {{.*}}>
+// CHECK:           %[[VAL_1:.*]] = llvm.getelementptr inbounds %[[VAL_0]][0, 2] : (!llvm.ptr) -> !llvm.ptr, !llvm.struct<"class.sycl::_V1::nd_item.3", {{.*}}>
 // CHECK:           %[[VAL_3:.*]] = llvm.getelementptr inbounds %[[VAL_1]][0, 3, 0, 0, 0] : (!llvm.ptr) -> !llvm.ptr, !llvm.struct<"class.sycl::_V1::group.3", {{.*}}>
 // CHECK:           %[[VAL_4:.*]] = llvm.load %[[VAL_3]] : !llvm.ptr -> i64
 // CHECK:           %[[VAL_5:.*]] = llvm.getelementptr inbounds %[[VAL_1]][0, 2, 0, 0, 1] : (!llvm.ptr) -> !llvm.ptr, !llvm.struct<"class.sycl::_V1::group.3", {{.*}}>