Fix after review

Signed-off-by: Zahira Ammarguellat <zahira.ammarguellat@intel.com>

Author: zahiraam

clang/lib/Sema/SemaSYCL.cpp

@@ -2372,15 +2372,6 @@ class SyclOptReportCreator : public SyclKernelFieldHandler {
   }
 
   bool handleSyclStreamType(FieldDecl *FD, QualType FieldTy) final {
-    // For the current implementation of stream class, the Visitor 'handles'
-    // stream argument and then visits each accessor field in stream. Therefore
-    // handleSpecialType in this case only adds a single argument for stream.
-    // The arguments corresponding to accessors in stream are handled in
-    // handleSyclAccessorType. The opt-report therefore does not diffrentiate
-    // between the accessors in streams and accessors captured by SYCL kernel.
-    // Once stream API is modified to use __init(), the visitor will no longer
-    // visit the stream object and opt-report output for stream class will be
-    // similar to that of other special types.
     return handleSpecialType(
         FD, FieldTy, KernelArgDescription(KernelArgDescription::Stream));
   }

clang/test/CodeGenSYCL/stream.cpp

@@ -1,18 +1,24 @@
 // RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o %t.ll
 // RUN: FileCheck < %t.ll --enable-var-scope %s
 //
-
 // CHECK: %[[RANGE_TYPE:"struct.*cl::sycl::range"]]
 // CHECK: %[[ID_TYPE:"struct.*cl::sycl::id"]]
+
 // CHECK: define dso_local spir_kernel void @{{.*}}StreamTester
 // CHECK-SAME: i8 addrspace(1)* [[ACC_DATA:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: %[[ID_TYPE]]* byval(%[[ID_TYPE]]) align 4 [[ACC_ID:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: i32 [[ACC_INT:%[a-zA-Z0-9_]+]])
 
-// CHECK: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::stream{{.*}} addrspace(4)* align 4 dereferenceable_or_null(16) %4, i8 addrspace(1)* %5, %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) {{.*}}%{{.*}}
-// CHECK: call spir_func void @{{.*}}__finalizeEv{{.*}}(%{{.*}}cl::sycl::stream{{.*}}" addrspace(4)* align 4 dereferenceable_or_null(16) %{{[0-9]+}})
+// Alloca and addrspace casts for kernel parameters
+// CHECK: [[ARG:%[a-zA-Z0-9_]+]].addr = alloca i8 addrspace(1)*, align 8
+// CHECK: [[ARG:%[a-zA-Z0-9_]+]].addr.ascast = addrspacecast i8 addrspace(1)** [[ARG]].addr to i8 addrspace(1)* addrspace(4)*
+// CHECK: [[ARG_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)* addrspace(4)* [[ARG]].addr.ascast, align 8,
+
+// Check __init and __finalize method calls
+// CHECK:  call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::stream" addrspace(4)* align 4 dereferenceable_or_null(16) %4, i8 addrspace(1)* [[ARG_LOAD]], %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) {{.*}}%{{.*}}
+// CHECK:  call spir_func void @_ZN2cl4sycl6stream10__finalizeEv(%{{.*}}cl::sycl::stream" addrspace(4)* align 4 dereferenceable_or_null(16) %{{[0-9]+}})
 
 #include "Inputs/sycl.hpp"
 

clang/test/SemaSYCL/streams.cpp

@@ -114,8 +114,9 @@ int main() {
 // element 1
 // CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 
-// HasArrayOfHasStreams struct
+// HasArrayOfHasStreams Array
 // CHECK: InitListExpr {{.*}} 'HasArrayOfHasStreams [2]'
+// // HasArrayOfHasStreams Struct
 // CHECK-NEXT: InitListExpr {{.*}} 'HasArrayOfHasStreams'
 // HasArrayOfHasStreams::i
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
@@ -142,6 +143,7 @@ int main() {
 // CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 // element 1
 // CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
+// HasArrayOfHasStreams Struct
 // CHECK-NEXT: InitListExpr {{.*}} 'HasArrayOfHasStreams'
 // HasArrayOfHasStreams::i
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
@@ -545,6 +547,7 @@ int main() {
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
+
 // _in_lambda_mdarray
 // [0][0]
 // CHECK: CXXMemberCallExpr {{.*}} 'void'