From 4490e85934da86510f929b83e99840f838504279 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Fri, 27 Sep 2024 13:50:39 +0800 Subject: [PATCH 01/11] [SYCLomatic] Use auto for template type Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/AnalysisInfo.cpp | 13 +++++++++---- clang/lib/DPCT/AnalysisInfo.h | 4 +++- 2 files changed, 12 insertions(+), 5 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index dbbb12ffbc93..b230dde97394 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -5224,9 +5224,11 @@ void DeviceFunctionInfo::mergeTextureObjectList( KernelCallExpr::ArgInfo::ArgInfo(const ParmVarDecl *PVD, KernelArgumentAnalysis &Analysis, const Expr *Arg, bool Used, int Index, - KernelCallExpr *BASE) + KernelCallExpr *BASE, const ParmVarDecl *TPVD) : IsPointer(false), IsRedeclareRequired(false), IsUsedAsLvalueAfterMalloc(Used), Index(Index) { + if (TPVD && TPVD->getType()->isDependentType()) + IsDependentType = true; if (isa(Arg)) { HasImplicitConversion = true; } else if (const auto *CCE = dyn_cast(Arg)) { @@ -5875,8 +5877,10 @@ void KernelCallExpr::buildArgsInfo(const CallExpr *CE) { if (auto *ArgDRE = dyn_cast(Arg->IgnoreImpCasts())) Used = isArgUsedAsLvalueUntil(ArgDRE, CE); const auto FD = CE->getDirectCallee(); - ArgsInfo.emplace_back(FD ? FD->parameters()[Idx] : nullptr, Analysis, Arg, - Used, Idx, this); + const FunctionTemplateDecl *FTD = FD ? FD->getPrimaryTemplate() : nullptr; + ArgsInfo.emplace_back( + FD ? FD->parameters()[Idx] : nullptr, Analysis, Arg, Used, Idx, this, + FTD ? FTD->getTemplatedDecl()->parameters()[Idx] : nullptr); } } } @@ -6226,7 +6230,8 @@ void KernelCallExpr::buildKernelArgsStmt() { } } else if (Arg.IsRedeclareRequired || IsInMacroDefine) { std::string TypeStr = "auto"; - if (Arg.HasImplicitConversion && !Arg.getTypeString().empty()) { + if (Arg.HasImplicitConversion && !Arg.getTypeString().empty() && + !Arg.IsDependentType) { TypeStr = Arg.getTypeString(); } SubmitStmts.CommandGroupList.emplace_back( diff --git a/clang/lib/DPCT/AnalysisInfo.h b/clang/lib/DPCT/AnalysisInfo.h index f0c718ea8410..87db36b8c9c2 100644 --- a/clang/lib/DPCT/AnalysisInfo.h +++ b/clang/lib/DPCT/AnalysisInfo.h @@ -2735,7 +2735,8 @@ class KernelCallExpr : public CallFunctionExpr { private: struct ArgInfo { ArgInfo(const ParmVarDecl *PVD, KernelArgumentAnalysis &Analysis, - const Expr *Arg, bool Used, int Index, KernelCallExpr *BASE); + const Expr *Arg, bool Used, int Index, KernelCallExpr *BASE, + const ParmVarDecl *TPVD); ArgInfo(const ParmVarDecl *PVD, const std::string &ArgsArrayName, KernelCallExpr *Kernel); ArgInfo(const ParmVarDecl *PVD, KernelCallExpr *Kernel); @@ -2762,6 +2763,7 @@ class KernelCallExpr : public CallFunctionExpr { bool IsDeviceRandomGeneratorType = false; bool HasImplicitConversion = false; bool IsDoublePointer = false; + bool IsDependentType = false; std::shared_ptr Texture; }; From c01f7f7f6b289ca4734c693843cc6c8c49536080 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Tue, 8 Oct 2024 11:15:26 +0800 Subject: [PATCH 02/11] Add test Signed-off-by: Jiang, Zhiwei --- clang/test/dpct/kernel-usm.cu | 27 +++++++++++++++++++++++++++ 1 file changed, 27 insertions(+) diff --git a/clang/test/dpct/kernel-usm.cu b/clang/test/dpct/kernel-usm.cu index 670d26d76583..f31628b68036 100644 --- a/clang/test/dpct/kernel-usm.cu +++ b/clang/test/dpct/kernel-usm.cu @@ -312,3 +312,30 @@ void run_foo13(float* a_host[]) { my_kernel5<<<1, 1>>>(a_host); } #endif + +template +__global__ void kernel6(const T* a); + +template +struct S { + T* getT() { return nullptr; } +}; + +// CHECK: void run_foo14(S s) { +// CHECK-NEXT: dpct::get_in_order_queue().submit( +// CHECK-NEXT: [&](sycl::handler &cgh) { +// CHECK-NEXT: auto s_getT_ct0 = s.getT(); +// CHECK-EMPTY: +// CHECK-NEXT: cgh.parallel_for>( +// CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), +// CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { +// CHECK-NEXT: kernel6(s_getT_ct0); +// CHECK-NEXT: }); +// CHECK-NEXT: }); +// CHECK-NEXT: } +void run_foo14(S s) { + kernel6<<<1, 1>>>(s.getT()); +} + +template +__global__ void kernel6(const T* a) {} From f4ca8b1a8089f39318da8262d6ee78080a7a4572 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Tue, 8 Oct 2024 15:36:04 +0800 Subject: [PATCH 03/11] Refine Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/AnalysisInfo.cpp | 9 +++- clang/test/dpct/decay.cu | 2 +- clang/test/dpct/group_local_memory.cu | 2 +- clang/test/dpct/kernel-call.cu | 16 ++++---- clang/test/dpct/kernel-usm.cu | 27 ------------ clang/test/dpct/kernel_without_name-usm.cu | 48 ++++++++++++++++++++++ clang/test/dpct/kernel_without_name.cu | 45 ++++++++++++++++++++ clang/test/dpct/sharedmem_var_dynamic.cu | 2 +- clang/test/dpct/sharedmem_var_static.cu | 2 +- clang/test/dpct/template-kernel-call.cu | 4 +- 10 files changed, 114 insertions(+), 43 deletions(-) create mode 100644 clang/test/dpct/kernel_without_name-usm.cu diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index b230dde97394..4730e913ca64 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -6199,6 +6199,9 @@ void KernelCallExpr::buildKernelArgsStmt() { if (Arg.IsDeviceRandomGeneratorType) { TypeStr = TypeStr + " *"; } + if (Arg.IsDependentType) { + TypeStr = "decltype(" + Arg.getArgString() + ")"; + } if (DpctGlobalInfo::isOptimizeMigration() && getFuncInfo() && !(getFuncInfo()->isParameterReferenced(ArgCounter))) { @@ -6230,10 +6233,12 @@ void KernelCallExpr::buildKernelArgsStmt() { } } else if (Arg.IsRedeclareRequired || IsInMacroDefine) { std::string TypeStr = "auto"; - if (Arg.HasImplicitConversion && !Arg.getTypeString().empty() && - !Arg.IsDependentType) { + if (Arg.HasImplicitConversion && !Arg.getTypeString().empty()) { TypeStr = Arg.getTypeString(); } + if (Arg.IsDependentType) { + TypeStr = "decltype(" + Arg.getArgString() + ")"; + } SubmitStmts.CommandGroupList.emplace_back( buildString(TypeStr, " ", Arg.getIdStringWithIndex(), " = ", Arg.getArgString(), ";")); diff --git a/clang/test/dpct/decay.cu b/clang/test/dpct/decay.cu index 5bd4aa00a80d..4ea352872422 100644 --- a/clang/test/dpct/decay.cu +++ b/clang/test/dpct/decay.cu @@ -16,7 +16,7 @@ __global__ void foo_kernel(T** R) //CHECK-NEXT: */ //CHECK-NEXT: dpct::get_out_of_order_queue().submit( //CHECK-NEXT: [&](sycl::handler &cgh) { -//CHECK-NEXT: dpct::access_wrapper R_acc_ct0(R, cgh); +//CHECK-NEXT: dpct::access_wrapper R_acc_ct0(R, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), diff --git a/clang/test/dpct/group_local_memory.cu b/clang/test/dpct/group_local_memory.cu index 6778239a1a47..0cf558c327ed 100644 --- a/clang/test/dpct/group_local_memory.cu +++ b/clang/test/dpct/group_local_memory.cu @@ -127,7 +127,7 @@ int main(void) { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, n), sycl::range<3>(1, 1, n)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: templateReverse((int *)(&d_d_acc_ct0[0]), n, item_ct1); + // CHECK-NEXT: templateReverse((decltype(d_d))(&d_d_acc_ct0[0]), n, item_ct1); // CHECK-NEXT: }); // CHECK-NEXT: }); templateReverse<<<1, n>>>(d_d, n); diff --git a/clang/test/dpct/kernel-call.cu b/clang/test/dpct/kernel-call.cu index 9764f08d7d39..f271e3bdb501 100644 --- a/clang/test/dpct/kernel-call.cu +++ b/clang/test/dpct/kernel-call.cu @@ -797,7 +797,7 @@ void run_foo13(float* a_host[]) { //CHECK-NEXT:dpct::get_out_of_order_queue().submit( //CHECK-NEXT: [&](sycl::handler &cgh) { //CHECK-NEXT: sycl::local_accessor aa_acc_ct1(cgh); - //CHECK-NEXT: dpct::access_wrapper a_host_acc_ct0(a_host, cgh); + //CHECK-NEXT: dpct::access_wrapper a_host_acc_ct0(a_host, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -834,8 +834,8 @@ void run_foo15() { double *d; //CHECK:q_ct1.submit( //CHECK-NEXT: [&](sycl::handler &cgh) { - //CHECK-NEXT: dpct::access_wrapper f_acc_ct0(f, cgh); - //CHECK-NEXT: dpct::access_wrapper f_acc_ct1(f, cgh); + //CHECK-NEXT: dpct::access_wrapper f_acc_ct0(f, cgh); + //CHECK-NEXT: dpct::access_wrapper f_acc_ct1(f, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -846,8 +846,8 @@ void run_foo15() { my_kernel7<<<1,1>>>(f, f); //CHECK:q_ct1.submit( //CHECK-NEXT: [&](sycl::handler &cgh) { - //CHECK-NEXT: dpct::access_wrapper d_acc_ct0(d, cgh); - //CHECK-NEXT: dpct::access_wrapper d_acc_ct1(d, cgh); + //CHECK-NEXT: dpct::access_wrapper d_acc_ct0(d, cgh); + //CHECK-NEXT: dpct::access_wrapper d_acc_ct1(d, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -867,7 +867,7 @@ void run_foo16() { float *fa, *fb; //CHECK:q_ct1.submit( //CHECK-NEXT: [&](sycl::handler &cgh) { - //CHECK-NEXT: dpct::access_wrapper fb_acc_ct1(fb, cgh); + //CHECK-NEXT: dpct::access_wrapper fb_acc_ct1(fb, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -878,12 +878,12 @@ void run_foo16() { my_kernel8<<<1,1>>>(fa, fb); //CHECK:q_ct1.submit( //CHECK-NEXT: [&](sycl::handler &cgh) { - //CHECK-NEXT: dpct::access_wrapper fb_acc_ct1(fb, cgh); + //CHECK-NEXT: dpct::access_wrapper fb_acc_ct1(fb, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), //CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - //CHECK-NEXT: my_kernel8((float *)nullptr, fb_acc_ct1.get_raw_pointer()); + //CHECK-NEXT: my_kernel8((decltype(fa))nullptr, fb_acc_ct1.get_raw_pointer()); //CHECK-NEXT: }); //CHECK-NEXT: }); my_kernel8<<<1,1>>>(fa, fb); diff --git a/clang/test/dpct/kernel-usm.cu b/clang/test/dpct/kernel-usm.cu index f31628b68036..670d26d76583 100644 --- a/clang/test/dpct/kernel-usm.cu +++ b/clang/test/dpct/kernel-usm.cu @@ -312,30 +312,3 @@ void run_foo13(float* a_host[]) { my_kernel5<<<1, 1>>>(a_host); } #endif - -template -__global__ void kernel6(const T* a); - -template -struct S { - T* getT() { return nullptr; } -}; - -// CHECK: void run_foo14(S s) { -// CHECK-NEXT: dpct::get_in_order_queue().submit( -// CHECK-NEXT: [&](sycl::handler &cgh) { -// CHECK-NEXT: auto s_getT_ct0 = s.getT(); -// CHECK-EMPTY: -// CHECK-NEXT: cgh.parallel_for>( -// CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), -// CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { -// CHECK-NEXT: kernel6(s_getT_ct0); -// CHECK-NEXT: }); -// CHECK-NEXT: }); -// CHECK-NEXT: } -void run_foo14(S s) { - kernel6<<<1, 1>>>(s.getT()); -} - -template -__global__ void kernel6(const T* a) {} diff --git a/clang/test/dpct/kernel_without_name-usm.cu b/clang/test/dpct/kernel_without_name-usm.cu new file mode 100644 index 000000000000..77f4981d93ef --- /dev/null +++ b/clang/test/dpct/kernel_without_name-usm.cu @@ -0,0 +1,48 @@ +// RUN: dpct --format-range=none -out-root %T/kernel_without_name-usm %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only +// RUN: FileCheck --input-file %T/kernel_without_name-usm/kernel_without_name-usm.dp.cpp --match-full-lines %s +// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/kernel_without_name-usm/kernel_without_name-usm.dp.cpp -o %T/kernel_without_name-usm/kernel_without_name-usm.dp.o %} + +template __global__ void foo_kernel1(const T *a); + +enum FLOATING_TYPE { FT_FLOAT, FT_DOUBLE }; + +struct Mat { + template U *data() { return (U *)_data; } + FLOATING_TYPE getType() { return _ft; } + + void *_data; + FLOATING_TYPE _ft; +}; + +#define DISPATCH(type, functor) \ + { \ + switch (type) { \ + case FT_FLOAT: { \ + using scalar_t = float; \ + functor(); \ + break; \ + } \ + case FT_DOUBLE: { \ + using scalar_t = double; \ + functor(); \ + break; \ + } \ + } \ + } + +void run_foo1(Mat mat) { + // CHECK: DISPATCH(mat.getType(), ([&] { dpct::get_in_order_queue().submit( + // CHECK-NEXT: [&](sycl::handler &cgh) { + // CHECK-NEXT: decltype(mat.data()) mat_data_scalar_t_ct0 = mat.data(); + // CHECK-EMPTY: + // CHECK-NEXT: cgh.parallel_for( + // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), + // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { + // CHECK-NEXT: foo_kernel1(mat_data_scalar_t_ct0); + // CHECK-NEXT: }); + // CHECK-NEXT: }); })); + DISPATCH(mat.getType(), ([&] { foo_kernel1<<<1, 1>>>(mat.data()); })); +} + +template __global__ void foo_kernel1(const T *a) {} +#undef DISPATCH diff --git a/clang/test/dpct/kernel_without_name.cu b/clang/test/dpct/kernel_without_name.cu index 6da5018876ed..ddb0590f38b0 100644 --- a/clang/test/dpct/kernel_without_name.cu +++ b/clang/test/dpct/kernel_without_name.cu @@ -468,3 +468,48 @@ void func() { // CHECK-NEXT: }); kfunc<<<128, 32>>>(der, 1, a2); } + +template __global__ void foo_kernel13(const T *a); + +enum FLOATING_TYPE { FT_FLOAT, FT_DOUBLE }; + +struct Mat { + template U *data() { return (U *)_data; } + FLOATING_TYPE getType() { return _ft; } + + void *_data; + FLOATING_TYPE _ft; +}; + +#define DISPATCH(type, functor) \ + { \ + switch (type) { \ + case FT_FLOAT: { \ + using scalar_t = float; \ + functor(); \ + break; \ + } \ + case FT_DOUBLE: { \ + using scalar_t = double; \ + functor(); \ + break; \ + } \ + } \ + } + +void run_foo13(Mat mat) { + // CHECK: DISPATCH(mat.getType(), ([&] { dpct::get_out_of_order_queue().submit( + // CHECK-NEXT: [&](sycl::handler &cgh) { + // CHECK-NEXT: dpct::access_wrapper())> mat_data_scalar_t_acc_ct0(mat.data(), cgh); + // CHECK-EMPTY: + // CHECK-NEXT: cgh.parallel_for( + // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), + // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { + // CHECK-NEXT: foo_kernel13(mat_data_scalar_t_acc_ct0.get_raw_pointer()); + // CHECK-NEXT: }); + // CHECK-NEXT: }); })); + DISPATCH(mat.getType(), ([&] { foo_kernel13<<<1, 1>>>(mat.data()); })); +} + +template __global__ void foo_kernel13(const T *a) {} +#undef DISPATCH diff --git a/clang/test/dpct/sharedmem_var_dynamic.cu b/clang/test/dpct/sharedmem_var_dynamic.cu index ffe0e6c8781a..023fc699b6ed 100644 --- a/clang/test/dpct/sharedmem_var_dynamic.cu +++ b/clang/test/dpct/sharedmem_var_dynamic.cu @@ -101,7 +101,7 @@ int main(void) { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, n), sycl::range<3>(1, 1, n)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: templateReverse((int *)(&d_d_acc_ct0[0]), n, item_ct1, dpct_local_acc_ct1.template get_multi_ptr().get()); + // CHECK-NEXT: templateReverse((decltype(d_d))(&d_d_acc_ct0[0]), n, item_ct1, dpct_local_acc_ct1.template get_multi_ptr().get()); // CHECK-NEXT: }); // CHECK-NEXT: }); templateReverse<<<1, n, 4>>>(d_d, n); diff --git a/clang/test/dpct/sharedmem_var_static.cu b/clang/test/dpct/sharedmem_var_static.cu index 1080480a75e1..7dd6e02de37d 100644 --- a/clang/test/dpct/sharedmem_var_static.cu +++ b/clang/test/dpct/sharedmem_var_static.cu @@ -189,7 +189,7 @@ int main(void) { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, n), sycl::range<3>(1, 1, n)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: templateReverse((int *)(&d_d_acc_ct0[0]), n, item_ct1, s_acc_ct1, s3_acc_ct1); + // CHECK-NEXT: templateReverse((decltype(d_d))(&d_d_acc_ct0[0]), n, item_ct1, s_acc_ct1, s3_acc_ct1); // CHECK-NEXT: }); // CHECK-NEXT: }); templateReverse<<<1, n>>>(d_d, n); diff --git a/clang/test/dpct/template-kernel-call.cu b/clang/test/dpct/template-kernel-call.cu index f598f458f650..9ae0e2979cb6 100644 --- a/clang/test/dpct/template-kernel-call.cu +++ b/clang/test/dpct/template-kernel-call.cu @@ -210,8 +210,8 @@ int main() { // CHECK-NEXT:*/ // CHECK-NEXT: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper karg1_acc_ct0((const LA *)karg1, cgh); - // CHECK-NEXT: dpct::access_wrapper karg2_acc_ct1(karg2, cgh); + // CHECK-NEXT: dpct::access_wrapper karg1_acc_ct0((const LA *)karg1, cgh); + // CHECK-NEXT: dpct::access_wrapper karg2_acc_ct1(karg2, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for, LA>>( // CHECK-NEXT: sycl::nd_range<3>(griddim * threaddim, threaddim), From 1fb6f001789e3c0d9bf95374238589b55b0855c3 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Wed, 9 Oct 2024 17:18:54 +0800 Subject: [PATCH 04/11] Update Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/AnalysisInfo.cpp | 12 +- clang/runtime/dpct-rt/include/dpct/memory.hpp | 16 +- clang/test/dpct/accessor-offset.cu | 48 ++--- clang/test/dpct/atomic_functions.cu | 2 +- clang/test/dpct/cpp_test.cpp | 4 +- clang/test/dpct/cuda_cache_config.cu | 2 +- clang/test/dpct/cuda_const.cu | 4 +- clang/test/dpct/cuda_kernel_include.cu | 2 +- clang/test/dpct/curand-device.cu | 10 +- clang/test/dpct/decay.cu | 2 +- clang/test/dpct/devicemem.cu | 4 +- clang/test/dpct/group_local_memory.cu | 10 +- .../dpct/kernel-call-inner-virtual-pointer.cu | 8 +- .../dpct/kernel-call-origcode-embedded.cu | 4 +- clang/test/dpct/kernel-call.cu | 46 ++--- clang/test/dpct/kernel-call_same_args.cu | 4 +- clang/test/dpct/kernel_without_name-usm.cu | 2 +- clang/test/dpct/kernel_without_name.cu | 16 +- clang/test/dpct/launch-kernel-cooperative.cu | 4 +- clang/test/dpct/launch-kernel.cu | 4 +- clang/test/dpct/mf-kernel.cu | 2 +- clang/test/dpct/multi-files-main.cu | 2 +- clang/test/dpct/sharedmem_var_dynamic.cu | 8 +- clang/test/dpct/sharedmem_var_static.cu | 8 +- clang/test/dpct/sycl_style_double2.cu | 2 +- clang/test/dpct/sycl_style_int2.cu | 2 +- clang/test/dpct/template-kernel-call.cu | 16 +- clang/test/dpct/template_initialization.cu | 4 +- clang/test/dpct/test_path_in_windows.cu | 4 +- clang/test/dpct/thrust-complex-usmnone.cu | 4 +- .../tm-nonusm-no-submit-barrier-profiling.cu | 14 +- clang/test/dpct/tm-nonusm-profiling.cu | 20 +- clang/test/dpct/vector_type.cu | 192 +++++++++--------- .../src/kernel_warp.cu | 4 +- clang/test/dpct/wildcard_test/abc.cu | 4 +- 35 files changed, 247 insertions(+), 243 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 4730e913ca64..23fc808c56d6 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -6215,7 +6215,7 @@ void KernelCallExpr::buildKernelArgsStmt() { if (Arg.IsUsedAsLvalueAfterMalloc) { requestFeature(HelperFeatureEnum::device_ext); SubmitStmts.AccessorList.emplace_back(buildString( - MapNames::getDpctNamespace() + "access_wrapper<", TypeStr, "> ", + MapNames::getDpctNamespace() + "access_wrapper ", Arg.getIdStringWithSuffix("acc"), "(", Arg.getArgString(), Arg.IsDefinedOnDevice ? ".get_ptr()" : "", ", cgh);")); KernelArgs += buildString(Arg.getIdStringWithSuffix("acc"), @@ -6227,18 +6227,16 @@ void KernelCallExpr::buildKernelArgsStmt() { " = " + MapNames::getDpctNamespace() + "get_access(", Arg.getArgString(), Arg.IsDefinedOnDevice ? ".get_ptr()" : "", ", cgh);")); - KernelArgs += buildString("(", TypeStr, ")(&", - Arg.getIdStringWithSuffix("acc"), "[0])"); + KernelArgs += + buildString("&", Arg.getIdStringWithSuffix("acc"), "[0]"); } } } else if (Arg.IsRedeclareRequired || IsInMacroDefine) { std::string TypeStr = "auto"; - if (Arg.HasImplicitConversion && !Arg.getTypeString().empty()) { + if (Arg.HasImplicitConversion && !Arg.getTypeString().empty() && + !Arg.IsDependentType) { TypeStr = Arg.getTypeString(); } - if (Arg.IsDependentType) { - TypeStr = "decltype(" + Arg.getArgString() + ")"; - } SubmitStmts.CommandGroupList.emplace_back( buildString(TypeStr, " ", Arg.getIdStringWithIndex(), " = ", Arg.getArgString(), ";")); diff --git a/clang/runtime/dpct-rt/include/dpct/memory.hpp b/clang/runtime/dpct-rt/include/dpct/memory.hpp index ce8aa699cc81..11328725f899 100644 --- a/clang/runtime/dpct-rt/include/dpct/memory.hpp +++ b/clang/runtime/dpct-rt/include/dpct/memory.hpp @@ -918,6 +918,8 @@ static buffer_t get_buffer(const void *ptr) { template class access_wrapper { + static_assert(!std::is_same_v, void>, + "dataT cannot be void"); sycl::accessor accessor; size_t offset; @@ -926,7 +928,7 @@ class access_wrapper { /// /// \param ptr Pointer to memory. /// \param cgh The command group handler. - access_wrapper(const void *ptr, sycl::handler &cgh) + access_wrapper(const dataT *ptr, sycl::handler &cgh) : accessor(get_buffer(ptr).get_access(cgh)), offset(0) { auto alloc = detail::mem_mgr::instance().translate_ptr(ptr); offset = (byte_t *)ptr - alloc.alloc_ptr; @@ -944,12 +946,16 @@ class access_wrapper { /// If NULL is passed as an argument, an exception will be thrown. /// \param cgh The command group handler. /// \returns an accessor. -template -static sycl::accessor -get_access(const void *ptr, sycl::handler &cgh) { +template +static auto get_access(const T *ptr, sycl::handler &cgh) { if (ptr) { auto alloc = detail::mem_mgr::instance().translate_ptr(ptr); - return alloc.buffer.get_access(cgh); + if (std::is_same_v, void>) + return alloc.buffer.get_access(cgh); + else + return alloc.buffer.reinterpret(sycl::range<1>(alloc.size / sizeof(T))) + .get_access(cgh); } else { throw std::runtime_error( "NULL pointer argument in get_access function is invalid"); diff --git a/clang/test/dpct/accessor-offset.cu b/clang/test/dpct/accessor-offset.cu index 3a9328fff94b..8a2b76a7242d 100644 --- a/clang/test/dpct/accessor-offset.cu +++ b/clang/test/dpct/accessor-offset.cu @@ -39,7 +39,7 @@ void foo() { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: hello((int *)(&d_a_acc_ct0[0]), 23); + // CHECK-NEXT: hello(&d_a_acc_ct0[0], 23); // CHECK-NEXT: }); // CHECK-NEXT: }); // CHECK-NEXT: } @@ -54,7 +54,7 @@ void foo() { // CHECK-NEXT: d_a += 2; // CHECK-NEXT: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a, cgh); + // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -75,7 +75,7 @@ void foo() { // CHECK-NEXT: d_a -= 2; // CHECK-NEXT: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a, cgh); + // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -96,7 +96,7 @@ void foo() { // CHECK-NEXT: d_a = d_a + 2; // CHECK-NEXT: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a, cgh); + // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -117,7 +117,7 @@ void foo() { // CHECK-NEXT: d_a = d_a - 2; // CHECK-NEXT: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a, cgh); + // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -138,7 +138,7 @@ void foo() { // CHECK-NEXT: *(&d_a + 1) = (int *)dpct::dpct_malloc(n * sizeof(float)); // CHECK-NEXT: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a, cgh); + // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -158,7 +158,7 @@ void foo() { // CHECK-NEXT: d_a = (int *)dpct::dpct_malloc(n * sizeof(float)); // CHECK-NEXT: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a + 1, cgh); + // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a + 1, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -178,7 +178,7 @@ void foo() { // CHECK-NEXT: *p = (int *)dpct::dpct_malloc(n * sizeof(float)); // CHECK-NEXT: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a, cgh); + // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -203,7 +203,7 @@ void foo() { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: hello((int *)(&d_a_acc_ct0[0]), 23); + // CHECK-NEXT: hello(&d_a_acc_ct0[0], 23); // CHECK-NEXT: }); // CHECK-NEXT: }); // CHECK-NEXT: } @@ -227,7 +227,7 @@ void foo() { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: hello((int *)(&d_a_acc_ct0[0]), 23); + // CHECK-NEXT: hello(&d_a_acc_ct0[0], 23); // CHECK-NEXT: }); // CHECK-NEXT: }); // CHECK-NEXT: } @@ -251,7 +251,7 @@ void foo() { // CHECK-NEXT: mod(&d_a); // CHECK-NEXT: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a, cgh); + // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -272,7 +272,7 @@ void foo() { // CHECK-NEXT: mod2(d_a); // CHECK-NEXT: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a, cgh); + // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -298,7 +298,7 @@ void foo() { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: hello((int *)(&d_a_acc_ct0[0]), 23); + // CHECK-NEXT: hello(&d_a_acc_ct0[0], 23); // CHECK-NEXT: }); // CHECK-NEXT: }); // CHECK-NEXT: } @@ -319,7 +319,7 @@ void foo() { // CHECK-NEXT: d_a = d_a - 2; // CHECK-NEXT: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a, cgh); + // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -356,7 +356,7 @@ void foo() { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: hello((int *)(&d_a_acc_ct0[0]), 23); + // CHECK-NEXT: hello(&d_a_acc_ct0[0], 23); // CHECK-NEXT: }); // CHECK-NEXT: }); // CHECK-NEXT: } @@ -383,7 +383,7 @@ void foo() { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: hello((int *)(&d_a_acc_ct0[0])); + // CHECK-NEXT: hello(&d_a_acc_ct0[0]); // CHECK-NEXT: }); // CHECK-NEXT: }); // CHECK-NEXT: if (n > 45) { @@ -394,7 +394,7 @@ void foo() { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: hello((int *)(&d_a_acc_ct0[0])); + // CHECK-NEXT: hello(&d_a_acc_ct0[0]); // CHECK-NEXT: }); // CHECK-NEXT: }); // CHECK-NEXT: } @@ -402,7 +402,7 @@ void foo() { // CHECK-NEXT: d_a += 2; // CHECK-NEXT: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a, cgh); + // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -438,13 +438,13 @@ void foo() { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: hello((int *)(&d_a_acc_ct0[0])); + // CHECK-NEXT: hello(&d_a_acc_ct0[0]); // CHECK-NEXT: }); // CHECK-NEXT: }); // CHECK-NEXT: } // CHECK-NEXT: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a, cgh); + // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -472,13 +472,13 @@ void foo() { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: hello((int *)(&d_a_acc_ct0[0]), 23); + // CHECK-NEXT: hello(&d_a_acc_ct0[0], 23); // CHECK-NEXT: }); // CHECK-NEXT: }); // CHECK-NEXT: } // CHECK-NEXT: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a, cgh); + // CHECK-NEXT: dpct::access_wrapper d_a_acc_ct0(d_a, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -501,7 +501,7 @@ void foo() { // CHECK-NEXT: d_a_global = (int *)dpct::dpct_malloc(n * sizeof(float)); // CHECK-NEXT: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper d_a_global_acc_ct0(d_a_global, cgh); + // CHECK-NEXT: dpct::access_wrapper d_a_global_acc_ct0(d_a_global, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -557,7 +557,7 @@ int testVectorAdd(void) // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, blocksPerGrid) * sycl::range<3>(1, 1, threadsPerBlock), sycl::range<3>(1, 1, threadsPerBlock)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: vectorAdd((const float *)(&d_A_acc_ct0[0]), (const float *)(&d_B_acc_ct1[0]), (float *)(&d_C_acc_ct2[0]), numElements); + // CHECK-NEXT: vectorAdd(&d_A_acc_ct0[0], &d_B_acc_ct1[0], &d_C_acc_ct2[0], numElements); // CHECK-NEXT: }); // CHECK-NEXT: }); vectorAdd<<>>(d_A, d_B, d_C, numElements); diff --git a/clang/test/dpct/atomic_functions.cu b/clang/test/dpct/atomic_functions.cu index 0d564b6afdd0..e376e48241fd 100644 --- a/clang/test/dpct/atomic_functions.cu +++ b/clang/test/dpct/atomic_functions.cu @@ -103,7 +103,7 @@ void InvokeKernel() { cudaMemcpy(dev_ptr, host.get(), size, cudaMemcpyHostToDevice); // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper dev_ptr_acc_ct0(dev_ptr, cgh); + // CHECK-NEXT: dpct::access_wrapper dev_ptr_acc_ct0(dev_ptr, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, k_threads_per_block), sycl::range<3>(1, 1, k_threads_per_block)), diff --git a/clang/test/dpct/cpp_test.cpp b/clang/test/dpct/cpp_test.cpp index ee4d830decc5..7017915b9b13 100644 --- a/clang/test/dpct/cpp_test.cpp +++ b/clang/test/dpct/cpp_test.cpp @@ -260,8 +260,8 @@ int kernel_test() { int karg3 = 80; // CHECK: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper karg1_acc_ct0((const int *)karg1, cgh); - // CHECK-NEXT: dpct::access_wrapper karg2_acc_ct1(karg2, cgh); + // CHECK-NEXT: dpct::access_wrapper karg1_acc_ct0((const int *)karg1, cgh); + // CHECK-NEXT: dpct::access_wrapper karg2_acc_ct1(karg2, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(griddim * threaddim, threaddim), diff --git a/clang/test/dpct/cuda_cache_config.cu b/clang/test/dpct/cuda_cache_config.cu index c89ddd580a27..3f802336152c 100644 --- a/clang/test/dpct/cuda_cache_config.cu +++ b/clang/test/dpct/cuda_cache_config.cu @@ -70,7 +70,7 @@ int main(int argc, char **argv) { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, size / 64) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: simple_kernel((float *)(&d_array_acc_ct0[0])); + // CHECK-NEXT: simple_kernel(&d_array_acc_ct0[0]); // CHECK-NEXT: }); // CHECK-NEXT: }); simple_kernel<<>>(d_array); diff --git a/clang/test/dpct/cuda_const.cu b/clang/test/dpct/cuda_const.cu index 06dbd26dbbb4..18890bd5af5c 100644 --- a/clang/test/dpct/cuda_const.cu +++ b/clang/test/dpct/cuda_const.cu @@ -165,7 +165,7 @@ int main(int argc, char **argv) { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, size / 64) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: simple_kernel((float *)(&d_array_acc_ct0[0]), item_ct1, const_angle_acc_ct1.get_multi_ptr().get(), const_ptr_acc_ct1.get_multi_ptr().get()); + // CHECK-NEXT: simple_kernel(&d_array_acc_ct0[0], item_ct1, const_angle_acc_ct1.get_multi_ptr().get(), const_ptr_acc_ct1.get_multi_ptr().get()); // CHECK-NEXT: }); // CHECK-NEXT: }); // CHECK-NEXT: } @@ -198,7 +198,7 @@ int main(int argc, char **argv) { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, size / 64) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: simple_kernel_one((float *)(&d_array_acc_ct0[0]), item_ct1, const_float_acc_ct1, const_one_acc_ct1); + // CHECK-NEXT: simple_kernel_one(&d_array_acc_ct0[0], item_ct1, const_float_acc_ct1, const_one_acc_ct1); // CHECK-NEXT: }); // CHECK-NEXT: }); // CHECK-NEXT: } diff --git a/clang/test/dpct/cuda_kernel_include.cu b/clang/test/dpct/cuda_kernel_include.cu index 6e2b29cdbc2c..2bff35639a0f 100644 --- a/clang/test/dpct/cuda_kernel_include.cu +++ b/clang/test/dpct/cuda_kernel_include.cu @@ -42,7 +42,7 @@ int main(int argc, char **argv) { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, size / 64) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: simple_kernel((float *)(&d_array_acc_ct0[0]), item_ct1); + // CHECK-NEXT: simple_kernel(&d_array_acc_ct0[0], item_ct1); // CHECK-NEXT: }); // CHECK-NEXT: }); simple_kernel<<>>(d_array); diff --git a/clang/test/dpct/curand-device.cu b/clang/test/dpct/curand-device.cu index 7e8f075dbd5a..ed258cd1d73d 100644 --- a/clang/test/dpct/curand-device.cu +++ b/clang/test/dpct/curand-device.cu @@ -78,7 +78,7 @@ int main(int argc, char **argv) { //CHECK-NEXT: DPCT1101:{{[0-9]+}}: 'WARP_SIZE' expression was replaced with a value. Modify the code to use the original expression, provided in comments, if it is correct. //CHECK-NEXT: */ //CHECK-NEXT: sycl::local_accessor counter_acc_ct1(sycl::range<1>(32/*WARP_SIZE*/), cgh); - //CHECK-NEXT: dpct::access_wrapper dOut_acc_ct0(dOut, cgh); + //CHECK-NEXT: dpct::access_wrapper dOut_acc_ct0(dOut, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, NBLOCKS) * sycl::range<3>(1, 1, WARP_SIZE), sycl::range<3>(1, 1, WARP_SIZE)), @@ -100,7 +100,7 @@ int main(int argc, char **argv) { //CHECK: q_ct1.submit( //CHECK-NEXT: [&](sycl::handler &cgh) { - //CHECK-NEXT: dpct::access_wrapper> *> RandomStates_acc_ct1(RandomStates, cgh); + //CHECK-NEXT: dpct::access_wrapper RandomStates_acc_ct1(RandomStates, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)), @@ -110,8 +110,8 @@ int main(int argc, char **argv) { //CHECK-NEXT: }); //CHECK-NEXT: q_ct1.submit( //CHECK-NEXT: [&](sycl::handler &cgh) { - //CHECK-NEXT: dpct::access_wrapper Image_acc_ct0(Image, cgh); - //CHECK-NEXT: dpct::access_wrapper> *> RandomStates_acc_ct1(RandomStates, cgh); + //CHECK-NEXT: dpct::access_wrapper Image_acc_ct0(Image, cgh); + //CHECK-NEXT: dpct::access_wrapper RandomStates_acc_ct1(RandomStates, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)), @@ -144,7 +144,7 @@ int foo() { //CHECK-NEXT: cgh.parallel_for( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)), //CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - //CHECK-NEXT: cuda_kernel_initRND(1234, (dpct::rng::device::rng_generator> *)(&RandomStates_acc_ct1[0]), item_ct1); + //CHECK-NEXT: cuda_kernel_initRND(1234, &RandomStates_acc_ct1[0], item_ct1); //CHECK-NEXT: }); //CHECK-NEXT: }); cuda_kernel_initRND<<<16,32>>>(1234, RandomStates); diff --git a/clang/test/dpct/decay.cu b/clang/test/dpct/decay.cu index 4ea352872422..5d04cfab4790 100644 --- a/clang/test/dpct/decay.cu +++ b/clang/test/dpct/decay.cu @@ -16,7 +16,7 @@ __global__ void foo_kernel(T** R) //CHECK-NEXT: */ //CHECK-NEXT: dpct::get_out_of_order_queue().submit( //CHECK-NEXT: [&](sycl::handler &cgh) { -//CHECK-NEXT: dpct::access_wrapper R_acc_ct0(R, cgh); +//CHECK-NEXT: dpct::access_wrapper R_acc_ct0(R, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), diff --git a/clang/test/dpct/devicemem.cu b/clang/test/dpct/devicemem.cu index c94bcf3285b6..395e1d8e8a61 100644 --- a/clang/test/dpct/devicemem.cu +++ b/clang/test/dpct/devicemem.cu @@ -109,7 +109,7 @@ int main() { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, threads_per_block), sycl::range<3>(1, 1, threads_per_block)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: kernel1((float *)(&d_out_acc_ct0[0]), item_ct1, in_acc_ct1.get_multi_ptr().get()); + // CHECK-NEXT: kernel1(&d_out_acc_ct0[0], item_ct1, in_acc_ct1.get_multi_ptr().get()); // CHECK-NEXT: }); // CHECK-NEXT: }); // CHECK-NEXT: } @@ -132,7 +132,7 @@ int main() { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, threads_per_block), sycl::range<3>(1, 1, threads_per_block)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: kernel2((float *)(&d_out_acc_ct0[0]), item_ct1, al_acc_ct1, fx_acc_ct1.get_multi_ptr().get(), fy_acc_ct1, tmp_acc_ct1.get_multi_ptr().get()); + // CHECK-NEXT: kernel2(&d_out_acc_ct0[0], item_ct1, al_acc_ct1, fx_acc_ct1.get_multi_ptr().get(), fy_acc_ct1, tmp_acc_ct1.get_multi_ptr().get()); // CHECK-NEXT: }); // CHECK-NEXT: }); // CHECK-NEXT: } diff --git a/clang/test/dpct/group_local_memory.cu b/clang/test/dpct/group_local_memory.cu index 0cf558c327ed..5bf92ddb3873 100644 --- a/clang/test/dpct/group_local_memory.cu +++ b/clang/test/dpct/group_local_memory.cu @@ -80,7 +80,7 @@ void testTemplate() { // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper d_d_acc_ct0(d_d, cgh); + // CHECK-NEXT: dpct::access_wrapper d_d_acc_ct0(d_d, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, n), sycl::range<3>(1, 1, n)), @@ -114,7 +114,7 @@ int main(void) { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, n), sycl::range<3>(1, 1, n)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: staticReverse((int *)(&d_d_acc_ct0[0]), n, item_ct1); + // CHECK-NEXT: staticReverse(&d_d_acc_ct0[0], n, item_ct1); // CHECK-NEXT: }); // CHECK-NEXT: }); staticReverse<<<1, n>>>(d_d, n); @@ -127,7 +127,7 @@ int main(void) { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, n), sycl::range<3>(1, 1, n)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: templateReverse((decltype(d_d))(&d_d_acc_ct0[0]), n, item_ct1); + // CHECK-NEXT: templateReverse(&d_d_acc_ct0[0], n, item_ct1); // CHECK-NEXT: }); // CHECK-NEXT: }); templateReverse<<<1, n>>>(d_d, n); @@ -139,7 +139,7 @@ int main(void) { // CHECK-NEXT: cgh.parallel_for>>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, n), sycl::range<3>(1, 1, n)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: nonTypeTemplateReverse((int *)(&d_d_acc_ct0[0]), n, item_ct1); + // CHECK-NEXT: nonTypeTemplateReverse(&d_d_acc_ct0[0], n, item_ct1); // CHECK-NEXT: }); // CHECK-NEXT: }); nonTypeTemplateReverse<<<1, n>>>(d_d, n); @@ -155,7 +155,7 @@ __global__ void foo(int *pd, int len) { smem[threadIdx.x] = 0; } // CHECK-NEXT: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { // CHECK-NEXT: sycl::local_accessor smem_acc_ct1(sycl::range<1>(shareSz), cgh); -// CHECK-NEXT: dpct::access_wrapper pd_acc_ct0(pd, cgh); +// CHECK-NEXT: dpct::access_wrapper pd_acc_ct0(pd, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 32) * sycl::range<3>(1, 1, 8), sycl::range<3>(1, 1, 8)), diff --git a/clang/test/dpct/kernel-call-inner-virtual-pointer.cu b/clang/test/dpct/kernel-call-inner-virtual-pointer.cu index b29ed829d59e..2b16c921fb3e 100644 --- a/clang/test/dpct/kernel-call-inner-virtual-pointer.cu +++ b/clang/test/dpct/kernel-call-inner-virtual-pointer.cu @@ -64,7 +64,7 @@ int main() { //CHECK-NEXT:*/ //CHECK-NEXT: q_ct1.submit( //CHECK-NEXT: [&](sycl::handler &cgh) { - //CHECK-NEXT: dpct::access_wrapper b1_acc_ct0(b1, cgh); + //CHECK-NEXT: dpct::access_wrapper b1_acc_ct0(b1, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -84,7 +84,7 @@ int main() { //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), //CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - //CHECK-NEXT: k2((AAA *)(&b2_acc_ct0[0])); + //CHECK-NEXT: k2(&b2_acc_ct0[0]); //CHECK-NEXT: }); //CHECK-NEXT: }); k2<<<1,1>>>(b2); @@ -94,7 +94,7 @@ int main() { //CHECK-NEXT:*/ //CHECK-NEXT: q_ct1.submit( //CHECK-NEXT: [&](sycl::handler &cgh) { - //CHECK-NEXT: dpct::access_wrapper a1_acc_ct0(a1, cgh); + //CHECK-NEXT: dpct::access_wrapper a1_acc_ct0(a1, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -109,7 +109,7 @@ int main() { //CHECK-NEXT:*/ //CHECK-NEXT: q_ct1.submit( //CHECK-NEXT: [&](sycl::handler &cgh) { - //CHECK-NEXT: dpct::access_wrapper a2_acc_ct0(a2, cgh); + //CHECK-NEXT: dpct::access_wrapper a2_acc_ct0(a2, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), diff --git a/clang/test/dpct/kernel-call-origcode-embedded.cu b/clang/test/dpct/kernel-call-origcode-embedded.cu index bdb3874e21e8..f3cf3230da1f 100644 --- a/clang/test/dpct/kernel-call-origcode-embedded.cu +++ b/clang/test/dpct/kernel-call-origcode-embedded.cu @@ -68,8 +68,8 @@ int main() { // CHECK-NEXT: karg2, karg3);*/ // CHECK: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper karg1_acc_ct0((const int *)karg1, cgh); - // CHECK-NEXT: dpct::access_wrapper karg2_acc_ct1(karg2, cgh); + // CHECK-NEXT: dpct::access_wrapper karg1_acc_ct0((const int *)karg1, cgh); + // CHECK-NEXT: dpct::access_wrapper karg2_acc_ct1(karg2, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(griddim * threaddim, threaddim), diff --git a/clang/test/dpct/kernel-call.cu b/clang/test/dpct/kernel-call.cu index f271e3bdb501..3fc419bf770b 100644 --- a/clang/test/dpct/kernel-call.cu +++ b/clang/test/dpct/kernel-call.cu @@ -135,8 +135,8 @@ int main() { // CHECK-NEXT: */ // CHECK-NEXT: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper karg1_acc_ct0((const int *)karg1, cgh); - // CHECK-NEXT: dpct::access_wrapper karg2_acc_ct1(karg2, cgh); + // CHECK-NEXT: dpct::access_wrapper karg1_acc_ct0((const int *)karg1, cgh); + // CHECK-NEXT: dpct::access_wrapper karg2_acc_ct1(karg2, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(griddim * threaddim, threaddim), @@ -172,8 +172,8 @@ int main() { // CHECK: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper args_arg1_acc_ct0(args.arg1, cgh); - // CHECK-NEXT: dpct::access_wrapper args_arg2_acc_ct1(args.arg2, cgh); + // CHECK-NEXT: dpct::access_wrapper args_arg1_acc_ct0(args.arg1, cgh); + // CHECK-NEXT: dpct::access_wrapper args_arg2_acc_ct1(args.arg2, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 2, 1), sycl::range<3>(1, 2, 1)), @@ -297,7 +297,7 @@ __global__ void foo_kernel3(int *d) { //CHECK-NEXT: if (1) //CHECK-NEXT: dpct::get_out_of_order_queue().submit( //CHECK-NEXT: [&](sycl::handler &cgh) { -//CHECK-NEXT: dpct::access_wrapper g_a_acc_ct0(&g_a[0], cgh); +//CHECK-NEXT: dpct::access_wrapper g_a_acc_ct0(&g_a[0], cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(c, sycl::range<3>(1, 1, 1)), @@ -319,7 +319,7 @@ void run_foo(dim3 c, dim3 d) { //CHECK-NEXT: */ //CHECK-NEXT: q_ct1.submit( //CHECK-NEXT: [&](sycl::handler &cgh) { -//CHECK-NEXT: dpct::access_wrapper g_a_acc_ct0(g_a, cgh); +//CHECK-NEXT: dpct::access_wrapper g_a_acc_ct0(g_a, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(c * d, d), @@ -330,7 +330,7 @@ void run_foo(dim3 c, dim3 d) { //CHECK-NEXT: else //CHECK-NEXT: q_ct1.submit( //CHECK-NEXT: [&](sycl::handler &cgh) { -//CHECK-NEXT: dpct::access_wrapper g_a_acc_ct0(g_a, cgh); +//CHECK-NEXT: dpct::access_wrapper g_a_acc_ct0(g_a, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(c, sycl::range<3>(1, 1, 1)), @@ -352,7 +352,7 @@ void run_foo2(dim3 c, dim3 d) { //CHECK-NEXT: */ //CHECK-NEXT: dpct::get_out_of_order_queue().submit( //CHECK-NEXT: [&](sycl::handler &cgh) { -//CHECK-NEXT: dpct::access_wrapper g_a_acc_ct0(g_a, cgh); +//CHECK-NEXT: dpct::access_wrapper g_a_acc_ct0(g_a, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(c * d, d), @@ -369,7 +369,7 @@ void run_foo3(dim3 c, dim3 d) { //CHECK-NEXT: while (1) //CHECK-NEXT: dpct::get_out_of_order_queue().submit( //CHECK-NEXT: [&](sycl::handler &cgh) { -//CHECK-NEXT: dpct::access_wrapper g_a_acc_ct0(g_a, cgh); +//CHECK-NEXT: dpct::access_wrapper g_a_acc_ct0(g_a, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(c, sycl::range<3>(1, 1, 1)), @@ -394,7 +394,7 @@ void run_foo4(dim3 c, dim3 d) { //CHECK-NEXT: dpct::get_out_of_order_queue().submit( //CHECK-NEXT: [&](sycl::handler &cgh) { //CHECK-NEXT: sycl::local_accessor resultInGroup_acc_ct1(sycl::range<1>(8), cgh); -//CHECK-NEXT: dpct::access_wrapper result_acc_ct0(result.get_ptr(), cgh); +//CHECK-NEXT: dpct::access_wrapper result_acc_ct0(result.get_ptr(), cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 4) * sycl::range<3>(1, 1, 8), sycl::range<3>(1, 1, 8)), @@ -420,7 +420,7 @@ int run_foo5 () { //CHECK-NEXT: dpct::get_out_of_order_queue().submit( //CHECK-NEXT: [&](sycl::handler &cgh) { //CHECK-NEXT: sycl::local_accessor resultInGroup_acc_ct1(sycl::range<1>(8), cgh); -//CHECK-NEXT: dpct::access_wrapper result2_acc_ct0(result2.get_ptr(), cgh); +//CHECK-NEXT: dpct::access_wrapper result2_acc_ct0(result2.get_ptr(), cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 4) * sycl::range<3>(1, 1, 8), sycl::range<3>(1, 1, 8)), @@ -441,7 +441,7 @@ int run_foo6 () { //CHECK-NEXT: dpct::get_out_of_order_queue().submit( //CHECK-NEXT: [&](sycl::handler &cgh) { //CHECK-NEXT: sycl::local_accessor resultInGroup_acc_ct1(sycl::range<1>(8), cgh); -//CHECK-NEXT: dpct::access_wrapper result3_acc_ct0(result3.get_ptr(), cgh); +//CHECK-NEXT: dpct::access_wrapper result3_acc_ct0(result3.get_ptr(), cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 4) * sycl::range<3>(1, 1, 8), sycl::range<3>(1, 1, 8)), @@ -468,7 +468,7 @@ int run_foo7 () { //CHECK-NEXT: in[0] = 42; //CHECK-NEXT: dpct::get_out_of_order_queue().submit( //CHECK-NEXT: [&](sycl::handler &cgh) { -//CHECK-NEXT: dpct::access_wrapper out_acc_ct1(out.get_ptr(), cgh); +//CHECK-NEXT: dpct::access_wrapper out_acc_ct1(out.get_ptr(), cgh); //CHECK-EMPTY: //CHECK-NEXT: auto in_ct0 = in[0]; //CHECK-EMPTY: @@ -544,7 +544,7 @@ __global__ void k(int *p){ //CHECK-NEXT: A aa; //CHECK-NEXT: q_ct1.submit( //CHECK-NEXT: [&](sycl::handler &cgh) { -//CHECK-NEXT: dpct::access_wrapper aa_get_pointer_acc_ct0(aa.get_pointer(), cgh); +//CHECK-NEXT: dpct::access_wrapper aa_get_pointer_acc_ct0(aa.get_pointer(), cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -554,7 +554,7 @@ __global__ void k(int *p){ //CHECK-NEXT: }); //CHECK-NEXT: q_ct1.submit( //CHECK-NEXT: [&](sycl::handler &cgh) { -//CHECK-NEXT: dpct::access_wrapper vec_get_pointer_acc_ct0(vec[2].get_pointer(), cgh); +//CHECK-NEXT: dpct::access_wrapper vec_get_pointer_acc_ct0(vec[2].get_pointer(), cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -770,7 +770,7 @@ int run_foo12() { static const int gg = i; //CHECK: dpct::get_out_of_order_queue().submit( //CHECK-NEXT: [&](sycl::handler &cgh) { - //CHECK-NEXT: dpct::access_wrapper bb_acc_ct1(bb, cgh); + //CHECK-NEXT: dpct::access_wrapper bb_acc_ct1(bb, cgh); //CHECK-EMPTY: //CHECK-NEXT: auto aa_ct0 = aa; //CHECK-NEXT: auto gg_ct6 = gg; @@ -797,7 +797,7 @@ void run_foo13(float* a_host[]) { //CHECK-NEXT:dpct::get_out_of_order_queue().submit( //CHECK-NEXT: [&](sycl::handler &cgh) { //CHECK-NEXT: sycl::local_accessor aa_acc_ct1(cgh); - //CHECK-NEXT: dpct::access_wrapper a_host_acc_ct0(a_host, cgh); + //CHECK-NEXT: dpct::access_wrapper a_host_acc_ct0(a_host, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -834,8 +834,8 @@ void run_foo15() { double *d; //CHECK:q_ct1.submit( //CHECK-NEXT: [&](sycl::handler &cgh) { - //CHECK-NEXT: dpct::access_wrapper f_acc_ct0(f, cgh); - //CHECK-NEXT: dpct::access_wrapper f_acc_ct1(f, cgh); + //CHECK-NEXT: dpct::access_wrapper f_acc_ct0(f, cgh); + //CHECK-NEXT: dpct::access_wrapper f_acc_ct1(f, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -846,8 +846,8 @@ void run_foo15() { my_kernel7<<<1,1>>>(f, f); //CHECK:q_ct1.submit( //CHECK-NEXT: [&](sycl::handler &cgh) { - //CHECK-NEXT: dpct::access_wrapper d_acc_ct0(d, cgh); - //CHECK-NEXT: dpct::access_wrapper d_acc_ct1(d, cgh); + //CHECK-NEXT: dpct::access_wrapper d_acc_ct0(d, cgh); + //CHECK-NEXT: dpct::access_wrapper d_acc_ct1(d, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -867,7 +867,7 @@ void run_foo16() { float *fa, *fb; //CHECK:q_ct1.submit( //CHECK-NEXT: [&](sycl::handler &cgh) { - //CHECK-NEXT: dpct::access_wrapper fb_acc_ct1(fb, cgh); + //CHECK-NEXT: dpct::access_wrapper fb_acc_ct1(fb, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -878,7 +878,7 @@ void run_foo16() { my_kernel8<<<1,1>>>(fa, fb); //CHECK:q_ct1.submit( //CHECK-NEXT: [&](sycl::handler &cgh) { - //CHECK-NEXT: dpct::access_wrapper fb_acc_ct1(fb, cgh); + //CHECK-NEXT: dpct::access_wrapper fb_acc_ct1(fb, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for>( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), diff --git a/clang/test/dpct/kernel-call_same_args.cu b/clang/test/dpct/kernel-call_same_args.cu index e4b54924c51a..ec0b80d674e3 100644 --- a/clang/test/dpct/kernel-call_same_args.cu +++ b/clang/test/dpct/kernel-call_same_args.cu @@ -17,8 +17,8 @@ int main() { int karg3 = 80; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper karg2_acc_ct0((const int *)karg2, cgh); - // CHECK-NEXT: dpct::access_wrapper karg2_acc_ct1(karg2, cgh); + // CHECK-NEXT: dpct::access_wrapper karg2_acc_ct0((const int *)karg2, cgh); + // CHECK-NEXT: dpct::access_wrapper karg2_acc_ct1(karg2, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(griddim * threaddim, threaddim), diff --git a/clang/test/dpct/kernel_without_name-usm.cu b/clang/test/dpct/kernel_without_name-usm.cu index 77f4981d93ef..3dc664bb976b 100644 --- a/clang/test/dpct/kernel_without_name-usm.cu +++ b/clang/test/dpct/kernel_without_name-usm.cu @@ -33,7 +33,7 @@ struct Mat { void run_foo1(Mat mat) { // CHECK: DISPATCH(mat.getType(), ([&] { dpct::get_in_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: decltype(mat.data()) mat_data_scalar_t_ct0 = mat.data(); + // CHECK-NEXT: auto mat_data_scalar_t_ct0 = mat.data(); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), diff --git a/clang/test/dpct/kernel_without_name.cu b/clang/test/dpct/kernel_without_name.cu index ddb0590f38b0..c469bc7b0a44 100644 --- a/clang/test/dpct/kernel_without_name.cu +++ b/clang/test/dpct/kernel_without_name.cu @@ -87,8 +87,8 @@ int main() { } args; //CHECK:q_ct1.submit( //CHECK-NEXT: [&](sycl::handler &cgh) { - //CHECK-NEXT: dpct::access_wrapper args_arg1_acc_ct0(args.arg1, cgh); - //CHECK-NEXT: dpct::access_wrapper args_arg2_acc_ct1(args.arg2, cgh); + //CHECK-NEXT: dpct::access_wrapper args_arg1_acc_ct0(args.arg1, cgh); + //CHECK-NEXT: dpct::access_wrapper args_arg2_acc_ct1(args.arg2, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for( testKernelPtr<<>>(args.arg1, args.arg2, karg3int); @@ -152,7 +152,7 @@ __global__ void foo_kernel3(int *d) { } //CHECK:dpct::get_out_of_order_queue().submit( //CHECK-NEXT: [&](sycl::handler &cgh) { -//CHECK-NEXT: dpct::access_wrapper g_a_acc_ct0(g_a, cgh); +//CHECK-NEXT: dpct::access_wrapper g_a_acc_ct0(g_a, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for( void run_foo(dim3 c, dim3 d) { @@ -165,12 +165,12 @@ void run_foo2(dim3 c, dim3 d) { //CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.out_of_order_queue(); //CHECK:q_ct1.submit( //CHECK-NEXT: [&](sycl::handler &cgh) { -//CHECK-NEXT: dpct::access_wrapper g_a_acc_ct0(g_a, cgh); +//CHECK-NEXT: dpct::access_wrapper g_a_acc_ct0(g_a, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for( //CHECK: q_ct1.submit( //CHECK-NEXT: [&](sycl::handler &cgh) { -//CHECK-NEXT: dpct::access_wrapper g_a_acc_ct0(g_a, cgh); +//CHECK-NEXT: dpct::access_wrapper g_a_acc_ct0(g_a, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for( if (1) @@ -180,7 +180,7 @@ void run_foo2(dim3 c, dim3 d) { } //CHECK:dpct::get_out_of_order_queue().submit( //CHECK-NEXT: [&](sycl::handler &cgh) { -//CHECK-NEXT: dpct::access_wrapper g_a_acc_ct0(g_a, cgh); +//CHECK-NEXT: dpct::access_wrapper g_a_acc_ct0(g_a, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for( void run_foo3(dim3 c, dim3 d) { @@ -189,7 +189,7 @@ void run_foo3(dim3 c, dim3 d) { } //CHECK:dpct::get_out_of_order_queue().submit( //CHECK-NEXT: [&](sycl::handler &cgh) { -//CHECK-NEXT: dpct::access_wrapper g_a_acc_ct0(g_a, cgh); +//CHECK-NEXT: dpct::access_wrapper g_a_acc_ct0(g_a, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for( void run_foo4(dim3 c, dim3 d) { @@ -500,7 +500,7 @@ struct Mat { void run_foo13(Mat mat) { // CHECK: DISPATCH(mat.getType(), ([&] { dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper())> mat_data_scalar_t_acc_ct0(mat.data(), cgh); + // CHECK-NEXT: dpct::access_wrapper mat_data_scalar_t_acc_ct0(mat.data(), cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), diff --git a/clang/test/dpct/launch-kernel-cooperative.cu b/clang/test/dpct/launch-kernel-cooperative.cu index 46a7f9bf83e4..4954f87307fe 100644 --- a/clang/test/dpct/launch-kernel-cooperative.cu +++ b/clang/test/dpct/launch-kernel-cooperative.cu @@ -48,7 +48,7 @@ int main() { // CHECK: static_cast *>(tex)->create_image(); // CHECK: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper d_acc_ct0(*(int **)args[0], cgh); + // CHECK-NEXT: dpct::access_wrapper d_acc_ct0(*(int **)args[0], cgh); // CHECK-EMPTY: // CHECK-NEXT: auto tex_acc = static_cast *>(*(dpct::image_wrapper_base_p *)args[1])->get_access(cgh); // CHECK-EMPTY: @@ -69,7 +69,7 @@ int main() { // CHECK-NEXT: [&](sycl::handler &cgh) { // CHECK-NEXT: sycl::local_accessor dpct_local_acc_ct1(sycl::range<1>(32), cgh); // CHECK-NEXT: sycl::local_accessor s_acc_ct1(sycl::range<1>(16), cgh); - // CHECK-NEXT: dpct::access_wrapper d_acc_ct0(*(int **)args[0], cgh); + // CHECK-NEXT: dpct::access_wrapper d_acc_ct0(*(int **)args[0], cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 16), sycl::range<3>(1, 1, 16)), diff --git a/clang/test/dpct/launch-kernel.cu b/clang/test/dpct/launch-kernel.cu index a29dd493cb18..416ee0d19fb6 100644 --- a/clang/test/dpct/launch-kernel.cu +++ b/clang/test/dpct/launch-kernel.cu @@ -46,7 +46,7 @@ int main() { // CHECK: static_cast *>(tex)->create_image(); // CHECK: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper d_acc_ct0(*(int **)args[0], cgh); + // CHECK-NEXT: dpct::access_wrapper d_acc_ct0(*(int **)args[0], cgh); // CHECK-EMPTY: // CHECK-NEXT: auto tex_acc = static_cast *>(*(dpct::image_wrapper_base_p *)args[1])->get_access(cgh); // CHECK-EMPTY: @@ -67,7 +67,7 @@ int main() { // CHECK-NEXT: [&](sycl::handler &cgh) { // CHECK-NEXT: sycl::local_accessor dpct_local_acc_ct1(sycl::range<1>(32), cgh); // CHECK-NEXT: sycl::local_accessor s_acc_ct1(sycl::range<1>(16), cgh); - // CHECK-NEXT: dpct::access_wrapper d_acc_ct0(*(int **)args[0], cgh); + // CHECK-NEXT: dpct::access_wrapper d_acc_ct0(*(int **)args[0], cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 16), sycl::range<3>(1, 1, 16)), diff --git a/clang/test/dpct/mf-kernel.cu b/clang/test/dpct/mf-kernel.cu index bce5f639f03f..e871e6dd04b2 100644 --- a/clang/test/dpct/mf-kernel.cu +++ b/clang/test/dpct/mf-kernel.cu @@ -75,7 +75,7 @@ __global__ void constAdd(float *C) { // CHECK-NEXT: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { // CHECK-NEXT: auto A_acc_ct1 = A_ct.get_access(cgh); -// CHECK-NEXT: dpct::access_wrapper d_C_acc_ct0(d_C, cgh); +// CHECK-NEXT: dpct::access_wrapper d_C_acc_ct0(d_C, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 3) * sycl::range<3>(1, 1, 3), sycl::range<3>(1, 1, 3)), diff --git a/clang/test/dpct/multi-files-main.cu b/clang/test/dpct/multi-files-main.cu index 0ec8a0df23a0..a59b7fa01b61 100644 --- a/clang/test/dpct/multi-files-main.cu +++ b/clang/test/dpct/multi-files-main.cu @@ -15,7 +15,7 @@ int main() { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 16), sycl::range<3>(1, 1, 16)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: simple_kernel((unsigned int *)(&i_array_acc_ct0[0]), item_ct1); + // CHECK-NEXT: simple_kernel(&i_array_acc_ct0[0], item_ct1); // CHECK-NEXT: }); // CHECK-NEXT: }); simple_kernel<<<16, 16>>>(i_array); diff --git a/clang/test/dpct/sharedmem_var_dynamic.cu b/clang/test/dpct/sharedmem_var_dynamic.cu index 023fc699b6ed..af458419f7ae 100644 --- a/clang/test/dpct/sharedmem_var_dynamic.cu +++ b/clang/test/dpct/sharedmem_var_dynamic.cu @@ -43,7 +43,7 @@ void testTemplate(){ // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { // CHECK-NEXT: sycl::local_accessor dpct_local_acc_ct1(sycl::range<1>(mem_size), cgh); - // CHECK-NEXT: dpct::access_wrapper d_d_acc_ct0(d_d, cgh); + // CHECK-NEXT: dpct::access_wrapper d_d_acc_ct0(d_d, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, n), sycl::range<3>(1, 1, n)), @@ -71,7 +71,7 @@ int main(void) { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, n), sycl::range<3>(1, 1, n)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: staticReverse((int *)(&d_d_acc_ct0[0]), n, item_ct1, dpct_local_acc_ct1.get_multi_ptr().get()); + // CHECK-NEXT: staticReverse(&d_d_acc_ct0[0], n, item_ct1, dpct_local_acc_ct1.get_multi_ptr().get()); // CHECK-NEXT: }); // CHECK-NEXT: }); staticReverse<<<1, n, mem_size>>>(d_d, n); @@ -88,7 +88,7 @@ int main(void) { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, n), sycl::range<3>(1, 1, n)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: staticReverse((int *)(&d_d_acc_ct0[0]), n, item_ct1, dpct_local_acc_ct1.get_multi_ptr().get()); + // CHECK-NEXT: staticReverse(&d_d_acc_ct0[0], n, item_ct1, dpct_local_acc_ct1.get_multi_ptr().get()); // CHECK-NEXT: }); // CHECK-NEXT: }); staticReverse<<<1, n, sizeof(int)>>>(d_d, n); @@ -101,7 +101,7 @@ int main(void) { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, n), sycl::range<3>(1, 1, n)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: templateReverse((decltype(d_d))(&d_d_acc_ct0[0]), n, item_ct1, dpct_local_acc_ct1.template get_multi_ptr().get()); + // CHECK-NEXT: templateReverse(&d_d_acc_ct0[0], n, item_ct1, dpct_local_acc_ct1.template get_multi_ptr().get()); // CHECK-NEXT: }); // CHECK-NEXT: }); templateReverse<<<1, n, 4>>>(d_d, n); diff --git a/clang/test/dpct/sharedmem_var_static.cu b/clang/test/dpct/sharedmem_var_static.cu index 7dd6e02de37d..3d719ab4e56e 100644 --- a/clang/test/dpct/sharedmem_var_static.cu +++ b/clang/test/dpct/sharedmem_var_static.cu @@ -101,7 +101,7 @@ void testTemplate() { // CHECK-NEXT: DPCT1101:{{[0-9]+}}: 'size' expression was replaced with a value. Modify the code to use the original expression, provided in comments, if it is correct. // CHECK-NEXT: */ // CHECK-NEXT: sycl::local_accessor s3_acc_ct1(cgh); - // CHECK-NEXT: dpct::access_wrapper d_d_acc_ct0(d_d, cgh); + // CHECK-NEXT: dpct::access_wrapper d_d_acc_ct0(d_d, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, n), sycl::range<3>(1, 1, n)), @@ -159,7 +159,7 @@ int main(void) { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, n), sycl::range<3>(1, 1, n)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: staticReverse((int *)(&d_d_acc_ct0[0]), n, item_ct1, a0_acc_ct1, s_acc_ct1.get_multi_ptr().get()); + // CHECK-NEXT: staticReverse(&d_d_acc_ct0[0], n, item_ct1, a0_acc_ct1, s_acc_ct1.get_multi_ptr().get()); // CHECK-NEXT: }); // CHECK-NEXT: }); staticReverse<<<1, n>>>(d_d, n); @@ -189,7 +189,7 @@ int main(void) { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, n), sycl::range<3>(1, 1, n)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: templateReverse((decltype(d_d))(&d_d_acc_ct0[0]), n, item_ct1, s_acc_ct1, s3_acc_ct1); + // CHECK-NEXT: templateReverse(&d_d_acc_ct0[0], n, item_ct1, s_acc_ct1, s3_acc_ct1); // CHECK-NEXT: }); // CHECK-NEXT: }); templateReverse<<<1, n>>>(d_d, n); @@ -202,7 +202,7 @@ int main(void) { // CHECK-NEXT: cgh.parallel_for>>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, n), sycl::range<3>(1, 1, n)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: nonTypeTemplateReverse((int *)(&d_d_acc_ct0[0]), n, item_ct1, s_acc_ct1.get_multi_ptr().get()); + // CHECK-NEXT: nonTypeTemplateReverse(&d_d_acc_ct0[0], n, item_ct1, s_acc_ct1.get_multi_ptr().get()); // CHECK-NEXT: }); // CHECK-NEXT: }); nonTypeTemplateReverse<<<1, n>>>(d_d, n); diff --git a/clang/test/dpct/sycl_style_double2.cu b/clang/test/dpct/sycl_style_double2.cu index a3e78f07cc29..1f2a3da86b9a 100644 --- a/clang/test/dpct/sycl_style_double2.cu +++ b/clang/test/dpct/sycl_style_double2.cu @@ -77,7 +77,7 @@ int main() { // CHECK: sycl::double2* data; // CHECK-NEXT: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper data_acc_ct0(data, cgh); + // CHECK-NEXT: dpct::access_wrapper data_acc_ct0(data, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), diff --git a/clang/test/dpct/sycl_style_int2.cu b/clang/test/dpct/sycl_style_int2.cu index ec9343f6525e..e368dcce05ba 100644 --- a/clang/test/dpct/sycl_style_int2.cu +++ b/clang/test/dpct/sycl_style_int2.cu @@ -91,7 +91,7 @@ int main() { // CHECK: sycl::int2* data; // CHECK-NEXT: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper data_acc_ct0(data, cgh); + // CHECK-NEXT: dpct::access_wrapper data_acc_ct0(data, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), diff --git a/clang/test/dpct/template-kernel-call.cu b/clang/test/dpct/template-kernel-call.cu index 9ae0e2979cb6..326a046413aa 100644 --- a/clang/test/dpct/template-kernel-call.cu +++ b/clang/test/dpct/template-kernel-call.cu @@ -98,8 +98,8 @@ void runTest() { // CHECK-NEXT:*/ // CHECK-NEXT: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper karg1_acc_ct0((const T *)karg1, cgh); - // CHECK-NEXT: dpct::access_wrapper karg2_acc_ct1(karg2, cgh); + // CHECK-NEXT: dpct::access_wrapper karg1_acc_ct0((const T *)karg1, cgh); + // CHECK-NEXT: dpct::access_wrapper karg2_acc_ct1(karg2, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for, T>>( // CHECK-NEXT: sycl::nd_range<3>(griddim * threaddim, threaddim), @@ -114,8 +114,8 @@ void runTest() { // CHECK-NEXT:*/ // CHECK-NEXT: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper karg1_acc_ct0((const T *)karg1, cgh); - // CHECK-NEXT: dpct::access_wrapper karg3_acc_ct1(karg3, cgh); + // CHECK-NEXT: dpct::access_wrapper karg1_acc_ct0((const T *)karg1, cgh); + // CHECK-NEXT: dpct::access_wrapper karg3_acc_ct1(karg3, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for, dpct_kernel_scalar, T>>( // CHECK-NEXT: sycl::nd_range<3>(griddim * threaddim, threaddim), @@ -130,8 +130,8 @@ void runTest() { // CHECK-NEXT:*/ // CHECK-NEXT: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper *> karg4_acc_ct0(karg4, cgh); - // CHECK-NEXT: dpct::access_wrapper karg5_acc_ct1(karg5, cgh); + // CHECK-NEXT: dpct::access_wrapper karg4_acc_ct0(karg4, cgh); + // CHECK-NEXT: dpct::access_wrapper karg5_acc_ct1(karg5, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for, TestTemplate>>( // CHECK-NEXT: sycl::nd_range<3>(griddim * threaddim, threaddim), @@ -210,8 +210,8 @@ int main() { // CHECK-NEXT:*/ // CHECK-NEXT: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper karg1_acc_ct0((const LA *)karg1, cgh); - // CHECK-NEXT: dpct::access_wrapper karg2_acc_ct1(karg2, cgh); + // CHECK-NEXT: dpct::access_wrapper karg1_acc_ct0((const LA *)karg1, cgh); + // CHECK-NEXT: dpct::access_wrapper karg2_acc_ct1(karg2, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for, LA>>( // CHECK-NEXT: sycl::nd_range<3>(griddim * threaddim, threaddim), diff --git a/clang/test/dpct/template_initialization.cu b/clang/test/dpct/template_initialization.cu index 3d2fa12efc35..f2d87452f355 100644 --- a/clang/test/dpct/template_initialization.cu +++ b/clang/test/dpct/template_initialization.cu @@ -52,13 +52,13 @@ void run_test() { // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper d_in_acc_ct0(d_in, cgh); + // CHECK-NEXT: dpct::access_wrapper d_in_acc_ct0(d_in, cgh); // CHECK-NEXT: auto d_out_acc_ct1 = dpct::get_access(d_out, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, num_threads), sycl::range<3>(1, 1, num_threads)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: kernel(d_in_acc_ct0.get_raw_pointer(), (T *)(&d_out_acc_ct1[0]), item_ct1); + // CHECK-NEXT: kernel(d_in_acc_ct0.get_raw_pointer(), &d_out_acc_ct1[0], item_ct1); // CHECK-NEXT: }); // CHECK-NEXT: }); kernel<<<1, num_threads>>>(d_in, d_out); diff --git a/clang/test/dpct/test_path_in_windows.cu b/clang/test/dpct/test_path_in_windows.cu index 8f81295ccb7e..2acf43831ee4 100644 --- a/clang/test/dpct/test_path_in_windows.cu +++ b/clang/test/dpct/test_path_in_windows.cu @@ -20,8 +20,8 @@ int main() { // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { // CHECK-NEXT: // accessors to device memory - // CHECK-NEXT: dpct::access_wrapper karg2_acc_ct0((const int *)karg2, cgh); - // CHECK-NEXT: dpct::access_wrapper karg2_acc_ct1(karg2, cgh); + // CHECK-NEXT: dpct::access_wrapper karg2_acc_ct0((const int *)karg2, cgh); + // CHECK-NEXT: dpct::access_wrapper karg2_acc_ct1(karg2, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(griddim * threaddim, threaddim), diff --git a/clang/test/dpct/thrust-complex-usmnone.cu b/clang/test/dpct/thrust-complex-usmnone.cu index ef8f83b195e0..187f7045fc03 100644 --- a/clang/test/dpct/thrust-complex-usmnone.cu +++ b/clang/test/dpct/thrust-complex-usmnone.cu @@ -106,8 +106,8 @@ int main() { // CHECK: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { // CHECK-NEXT: sycl::local_accessor, 1> s_acc_ct1(sycl::range<1>(10), cgh); - // CHECK-NEXT: dpct::access_wrapper *> cdp_acc_ct0(reinterpret_cast *>(cdp), cgh); - // CHECK-NEXT: dpct::access_wrapper *> thrust_raw_pointer_cast_dc_ptr_acc_ct2(dpct::get_raw_pointer(dc_ptr), cgh); + // CHECK-NEXT: dpct::access_wrapper cdp_acc_ct0(reinterpret_cast *>(cdp), cgh); + // CHECK-NEXT: dpct::access_wrapper thrust_raw_pointer_cast_dc_ptr_acc_ct2(dpct::get_raw_pointer(dc_ptr), cgh); // CHECK-EMPTY: // CHECK-NEXT: std::complex static_cast_thrust_complex_double_cdp_ct1 = static_cast>(*cdp); // CHECK-EMPTY: diff --git a/clang/test/dpct/tm-nonusm-no-submit-barrier-profiling.cu b/clang/test/dpct/tm-nonusm-no-submit-barrier-profiling.cu index e1c14d8398ac..f035799b0288 100644 --- a/clang/test/dpct/tm-nonusm-no-submit-barrier-profiling.cu +++ b/clang/test/dpct/tm-nonusm-no-submit-barrier-profiling.cu @@ -522,8 +522,8 @@ void RunTest() { // CHECK: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { -// CHECK-NEXT: dpct::access_wrapper d_idata_acc_ct0(d_idata, cgh); -// CHECK-NEXT: dpct::access_wrapper d_block_sums_acc_ct1(d_block_sums, cgh); +// CHECK-NEXT: dpct::access_wrapper d_idata_acc_ct0(d_idata, cgh); +// CHECK-NEXT: dpct::access_wrapper d_block_sums_acc_ct1(d_block_sums, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, num_threads), sycl::range<3>(1, 1, num_threads)), @@ -588,9 +588,9 @@ void test_1999(void* ref_image, void* cur_image, // CHECK: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { -// CHECK-NEXT: dpct::access_wrapper d_sads_acc_ct0(d_sads, cgh); -// CHECK-NEXT: dpct::access_wrapper d_cur_image_acc_ct1(d_cur_image, cgh); -// CHECK-NEXT: dpct::access_wrapper imgRef_acc_ct4(imgRef, cgh); +// CHECK-NEXT: dpct::access_wrapper d_sads_acc_ct0(d_sads, cgh); +// CHECK-NEXT: dpct::access_wrapper d_cur_image_acc_ct1(d_cur_image, cgh); +// CHECK-NEXT: dpct::access_wrapper imgRef_acc_ct4(imgRef, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(foo_kernel_1_blocks_in_grid * foo_kernel_1_threads_in_block, foo_kernel_1_threads_in_block), @@ -623,7 +623,7 @@ void test_1999(void* ref_image, void* cur_image, // CHECK: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { -// CHECK-NEXT: dpct::access_wrapper d_sads_acc_ct0(d_sads, cgh); +// CHECK-NEXT: dpct::access_wrapper d_sads_acc_ct0(d_sads, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(foo_kernel_2_blocks_in_grid * foo_kernel_2_threads_in_block, foo_kernel_2_threads_in_block), @@ -655,7 +655,7 @@ void test_1999(void* ref_image, void* cur_image, // CHECK: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { -// CHECK-NEXT: dpct::access_wrapper d_sads_acc_ct0(d_sads, cgh); +// CHECK-NEXT: dpct::access_wrapper d_sads_acc_ct0(d_sads, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(foo_kernel_3_blocks_in_grid * foo_kernel_3_threads_in_block, foo_kernel_3_threads_in_block), diff --git a/clang/test/dpct/tm-nonusm-profiling.cu b/clang/test/dpct/tm-nonusm-profiling.cu index 5b94e4e2c58d..dd896f0bd332 100644 --- a/clang/test/dpct/tm-nonusm-profiling.cu +++ b/clang/test/dpct/tm-nonusm-profiling.cu @@ -331,7 +331,7 @@ void foo() // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(gridSize * blockSize, blockSize), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { -// CHECK-NEXT: readTexels(kernelRepFactor, (float *)(&d_out_acc_ct1[0]), width); +// CHECK-NEXT: readTexels(kernelRepFactor, &d_out_acc_ct1[0], width); // CHECK-NEXT: }); // CHECK-NEXT: }); readTexels<<>>(kernelRepFactor, d_out, @@ -364,7 +364,7 @@ void foo() // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(gridSize * blockSize, blockSize), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { -// CHECK-NEXT: readTexelsFoo1(kernelRepFactor, (float *)(&d_out_acc_ct1[0])); +// CHECK-NEXT: readTexelsFoo1(kernelRepFactor, &d_out_acc_ct1[0]); // CHECK-NEXT: }); // CHECK-NEXT: }); readTexelsFoo1<<>> @@ -398,7 +398,7 @@ void foo() // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(gridSize * blockSize, blockSize), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { -// CHECK-NEXT: readTexelsFoo2(kernelRepFactor, (float *)(&d_out_acc_ct1[0]), width, height); +// CHECK-NEXT: readTexelsFoo2(kernelRepFactor, &d_out_acc_ct1[0], width, height); // CHECK-NEXT: }); // CHECK-NEXT: }); readTexelsFoo2<<>> @@ -536,8 +536,8 @@ void RunTest() { // CHECK: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { -// CHECK-NEXT: dpct::access_wrapper d_idata_acc_ct0(d_idata, cgh); -// CHECK-NEXT: dpct::access_wrapper d_block_sums_acc_ct1(d_block_sums, cgh); +// CHECK-NEXT: dpct::access_wrapper d_idata_acc_ct0(d_idata, cgh); +// CHECK-NEXT: dpct::access_wrapper d_block_sums_acc_ct1(d_block_sums, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, num_threads), sycl::range<3>(1, 1, num_threads)), @@ -600,9 +600,9 @@ void test_1999(void* ref_image, void* cur_image, // CHECK: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { -// CHECK-NEXT: dpct::access_wrapper d_sads_acc_ct0(d_sads, cgh); -// CHECK-NEXT: dpct::access_wrapper d_cur_image_acc_ct1(d_cur_image, cgh); -// CHECK-NEXT: dpct::access_wrapper imgRef_acc_ct4(imgRef, cgh); +// CHECK-NEXT: dpct::access_wrapper d_sads_acc_ct0(d_sads, cgh); +// CHECK-NEXT: dpct::access_wrapper d_cur_image_acc_ct1(d_cur_image, cgh); +// CHECK-NEXT: dpct::access_wrapper imgRef_acc_ct4(imgRef, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(foo_kernel_1_blocks_in_grid * foo_kernel_1_threads_in_block, foo_kernel_1_threads_in_block), @@ -631,7 +631,7 @@ void test_1999(void* ref_image, void* cur_image, // CHECK: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { -// CHECK-NEXT: dpct::access_wrapper d_sads_acc_ct0(d_sads, cgh); +// CHECK-NEXT: dpct::access_wrapper d_sads_acc_ct0(d_sads, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(foo_kernel_2_blocks_in_grid * foo_kernel_2_threads_in_block, foo_kernel_2_threads_in_block), @@ -658,7 +658,7 @@ void test_1999(void* ref_image, void* cur_image, // CHECK: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { -// CHECK-NEXT: dpct::access_wrapper d_sads_acc_ct0(d_sads, cgh); +// CHECK-NEXT: dpct::access_wrapper d_sads_acc_ct0(d_sads, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(foo_kernel_3_blocks_in_grid * foo_kernel_3_threads_in_block, foo_kernel_3_threads_in_block), diff --git a/clang/test/dpct/vector_type.cu b/clang/test/dpct/vector_type.cu index aa86aa9e3619..4512aca289a8 100644 --- a/clang/test/dpct/vector_type.cu +++ b/clang/test/dpct/vector_type.cu @@ -59,8 +59,8 @@ int main_char1() { int *char1_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper char1_e_acc_ct0(char1_e, cgh); - // CHECK-NEXT: dpct::access_wrapper char1_cast_acc_ct1((int8_t *)char1_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper char1_e_acc_ct0(char1_e, cgh); + // CHECK-NEXT: dpct::access_wrapper char1_cast_acc_ct1((int8_t *)char1_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -131,8 +131,8 @@ int main_char2() { int *char2_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper char2_e_acc_ct0(char2_e, cgh); - // CHECK-NEXT: dpct::access_wrapper char2_cast_acc_ct1((sycl::char2 *)char2_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper char2_e_acc_ct0(char2_e, cgh); + // CHECK-NEXT: dpct::access_wrapper char2_cast_acc_ct1((sycl::char2 *)char2_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -203,8 +203,8 @@ int main_char3() { int *char3_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper char3_e_acc_ct0(char3_e, cgh); - // CHECK-NEXT: dpct::access_wrapper char3_cast_acc_ct1((sycl::char3 *)char3_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper char3_e_acc_ct0(char3_e, cgh); + // CHECK-NEXT: dpct::access_wrapper char3_cast_acc_ct1((sycl::char3 *)char3_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -275,8 +275,8 @@ int main_char4() { int *char4_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper char4_e_acc_ct0(char4_e, cgh); - // CHECK-NEXT: dpct::access_wrapper char4_cast_acc_ct1((sycl::char4 *)char4_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper char4_e_acc_ct0(char4_e, cgh); + // CHECK-NEXT: dpct::access_wrapper char4_cast_acc_ct1((sycl::char4 *)char4_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -347,8 +347,8 @@ int main_double1() { int *double1_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper double1_e_acc_ct0(double1_e, cgh); - // CHECK-NEXT: dpct::access_wrapper double1_cast_acc_ct1((double *)double1_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper double1_e_acc_ct0(double1_e, cgh); + // CHECK-NEXT: dpct::access_wrapper double1_cast_acc_ct1((double *)double1_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -419,8 +419,8 @@ int main_double2() { int *double2_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper double2_e_acc_ct0(double2_e, cgh); - // CHECK-NEXT: dpct::access_wrapper double2_cast_acc_ct1((sycl::double2 *)double2_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper double2_e_acc_ct0(double2_e, cgh); + // CHECK-NEXT: dpct::access_wrapper double2_cast_acc_ct1((sycl::double2 *)double2_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -491,8 +491,8 @@ int main_double3() { int *double3_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper double3_e_acc_ct0(double3_e, cgh); - // CHECK-NEXT: dpct::access_wrapper double3_cast_acc_ct1((sycl::double3 *)double3_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper double3_e_acc_ct0(double3_e, cgh); + // CHECK-NEXT: dpct::access_wrapper double3_cast_acc_ct1((sycl::double3 *)double3_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -563,8 +563,8 @@ int main_double4() { int *double4_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper double4_e_acc_ct0(double4_e, cgh); - // CHECK-NEXT: dpct::access_wrapper double4_cast_acc_ct1((sycl::double4 *)double4_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper double4_e_acc_ct0(double4_e, cgh); + // CHECK-NEXT: dpct::access_wrapper double4_cast_acc_ct1((sycl::double4 *)double4_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -635,8 +635,8 @@ int main_float1() { int *float1_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper float1_e_acc_ct0(float1_e, cgh); - // CHECK-NEXT: dpct::access_wrapper float1_cast_acc_ct1((float *)float1_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper float1_e_acc_ct0(float1_e, cgh); + // CHECK-NEXT: dpct::access_wrapper float1_cast_acc_ct1((float *)float1_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -707,8 +707,8 @@ int main_float2() { int *float2_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper float2_e_acc_ct0(float2_e, cgh); - // CHECK-NEXT: dpct::access_wrapper float2_cast_acc_ct1((sycl::float2 *)float2_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper float2_e_acc_ct0(float2_e, cgh); + // CHECK-NEXT: dpct::access_wrapper float2_cast_acc_ct1((sycl::float2 *)float2_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -779,8 +779,8 @@ int main_float3() { int *float3_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper float3_e_acc_ct0(float3_e, cgh); - // CHECK-NEXT: dpct::access_wrapper float3_cast_acc_ct1((sycl::float3 *)float3_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper float3_e_acc_ct0(float3_e, cgh); + // CHECK-NEXT: dpct::access_wrapper float3_cast_acc_ct1((sycl::float3 *)float3_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -851,8 +851,8 @@ int main_float4() { int *float4_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper float4_e_acc_ct0(float4_e, cgh); - // CHECK-NEXT: dpct::access_wrapper float4_cast_acc_ct1((sycl::float4 *)float4_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper float4_e_acc_ct0(float4_e, cgh); + // CHECK-NEXT: dpct::access_wrapper float4_cast_acc_ct1((sycl::float4 *)float4_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -923,8 +923,8 @@ int main_int1() { int *int1_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper int1_e_acc_ct0(int1_e, cgh); - // CHECK-NEXT: dpct::access_wrapper int1_cast_acc_ct1((int32_t *)int1_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper int1_e_acc_ct0(int1_e, cgh); + // CHECK-NEXT: dpct::access_wrapper int1_cast_acc_ct1((int32_t *)int1_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -995,8 +995,8 @@ int main_int2() { int *int2_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper int2_e_acc_ct0(int2_e, cgh); - // CHECK-NEXT: dpct::access_wrapper int2_cast_acc_ct1((sycl::int2 *)int2_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper int2_e_acc_ct0(int2_e, cgh); + // CHECK-NEXT: dpct::access_wrapper int2_cast_acc_ct1((sycl::int2 *)int2_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -1067,8 +1067,8 @@ int main_int3() { int *int3_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper int3_e_acc_ct0(int3_e, cgh); - // CHECK-NEXT: dpct::access_wrapper int3_cast_acc_ct1((sycl::int3 *)int3_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper int3_e_acc_ct0(int3_e, cgh); + // CHECK-NEXT: dpct::access_wrapper int3_cast_acc_ct1((sycl::int3 *)int3_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -1141,8 +1141,8 @@ int main_int4() { int *int4_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper int4_e_acc_ct0(int4_e, cgh); - // CHECK-NEXT: dpct::access_wrapper int4_cast_acc_ct1((sycl::int4 *)int4_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper int4_e_acc_ct0(int4_e, cgh); + // CHECK-NEXT: dpct::access_wrapper int4_cast_acc_ct1((sycl::int4 *)int4_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -1213,8 +1213,8 @@ int main_long1() { int *long1_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper long1_e_acc_ct0(long1_e, cgh); - // CHECK-NEXT: dpct::access_wrapper long1_cast_acc_ct1((int64_t *)long1_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper long1_e_acc_ct0(long1_e, cgh); + // CHECK-NEXT: dpct::access_wrapper long1_cast_acc_ct1((int64_t *)long1_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -1285,8 +1285,8 @@ int main_long2() { int *long2_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper long2_e_acc_ct0(long2_e, cgh); - // CHECK-NEXT: dpct::access_wrapper long2_cast_acc_ct1((sycl::long2 *)long2_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper long2_e_acc_ct0(long2_e, cgh); + // CHECK-NEXT: dpct::access_wrapper long2_cast_acc_ct1((sycl::long2 *)long2_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -1357,8 +1357,8 @@ int main_long3() { int *long3_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper long3_e_acc_ct0(long3_e, cgh); - // CHECK-NEXT: dpct::access_wrapper long3_cast_acc_ct1((sycl::long3 *)long3_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper long3_e_acc_ct0(long3_e, cgh); + // CHECK-NEXT: dpct::access_wrapper long3_cast_acc_ct1((sycl::long3 *)long3_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -1429,8 +1429,8 @@ int main_long4() { int *long4_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper long4_e_acc_ct0(long4_e, cgh); - // CHECK-NEXT: dpct::access_wrapper long4_cast_acc_ct1((sycl::long4 *)long4_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper long4_e_acc_ct0(long4_e, cgh); + // CHECK-NEXT: dpct::access_wrapper long4_cast_acc_ct1((sycl::long4 *)long4_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -1501,8 +1501,8 @@ int main_longlong1() { int *longlong1_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper longlong1_e_acc_ct0(longlong1_e, cgh); - // CHECK-NEXT: dpct::access_wrapper longlong1_cast_acc_ct1((int64_t *)longlong1_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper longlong1_e_acc_ct0(longlong1_e, cgh); + // CHECK-NEXT: dpct::access_wrapper longlong1_cast_acc_ct1((int64_t *)longlong1_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -1573,8 +1573,8 @@ int main_longlong2() { int *longlong2_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper longlong2_e_acc_ct0(longlong2_e, cgh); - // CHECK-NEXT: dpct::access_wrapper longlong2_cast_acc_ct1((sycl::long2 *)longlong2_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper longlong2_e_acc_ct0(longlong2_e, cgh); + // CHECK-NEXT: dpct::access_wrapper longlong2_cast_acc_ct1((sycl::long2 *)longlong2_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -1645,8 +1645,8 @@ int main_longlong3() { int *longlong3_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper longlong3_e_acc_ct0(longlong3_e, cgh); - // CHECK-NEXT: dpct::access_wrapper longlong3_cast_acc_ct1((sycl::long3 *)longlong3_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper longlong3_e_acc_ct0(longlong3_e, cgh); + // CHECK-NEXT: dpct::access_wrapper longlong3_cast_acc_ct1((sycl::long3 *)longlong3_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -1717,8 +1717,8 @@ int main_longlong4() { int *longlong4_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper longlong4_e_acc_ct0(longlong4_e, cgh); - // CHECK-NEXT: dpct::access_wrapper longlong4_cast_acc_ct1((sycl::long4 *)longlong4_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper longlong4_e_acc_ct0(longlong4_e, cgh); + // CHECK-NEXT: dpct::access_wrapper longlong4_cast_acc_ct1((sycl::long4 *)longlong4_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -1789,8 +1789,8 @@ int main_short1() { int *short1_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper short1_e_acc_ct0(short1_e, cgh); - // CHECK-NEXT: dpct::access_wrapper short1_cast_acc_ct1((int16_t *)short1_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper short1_e_acc_ct0(short1_e, cgh); + // CHECK-NEXT: dpct::access_wrapper short1_cast_acc_ct1((int16_t *)short1_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -1861,8 +1861,8 @@ int main_short2() { int *short2_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper short2_e_acc_ct0(short2_e, cgh); - // CHECK-NEXT: dpct::access_wrapper short2_cast_acc_ct1((sycl::short2 *)short2_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper short2_e_acc_ct0(short2_e, cgh); + // CHECK-NEXT: dpct::access_wrapper short2_cast_acc_ct1((sycl::short2 *)short2_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -1933,8 +1933,8 @@ int main_short3() { int *short3_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper short3_e_acc_ct0(short3_e, cgh); - // CHECK-NEXT: dpct::access_wrapper short3_cast_acc_ct1((sycl::short3 *)short3_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper short3_e_acc_ct0(short3_e, cgh); + // CHECK-NEXT: dpct::access_wrapper short3_cast_acc_ct1((sycl::short3 *)short3_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -2005,8 +2005,8 @@ int main_short4() { int *short4_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper short4_e_acc_ct0(short4_e, cgh); - // CHECK-NEXT: dpct::access_wrapper short4_cast_acc_ct1((sycl::short4 *)short4_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper short4_e_acc_ct0(short4_e, cgh); + // CHECK-NEXT: dpct::access_wrapper short4_cast_acc_ct1((sycl::short4 *)short4_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -2077,8 +2077,8 @@ int main_uchar1() { int *uchar1_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper uchar1_e_acc_ct0(uchar1_e, cgh); - // CHECK-NEXT: dpct::access_wrapper uchar1_cast_acc_ct1((uint8_t *)uchar1_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper uchar1_e_acc_ct0(uchar1_e, cgh); + // CHECK-NEXT: dpct::access_wrapper uchar1_cast_acc_ct1((uint8_t *)uchar1_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -2149,8 +2149,8 @@ int main_uchar2() { int *uchar2_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper uchar2_e_acc_ct0(uchar2_e, cgh); - // CHECK-NEXT: dpct::access_wrapper uchar2_cast_acc_ct1((sycl::uchar2 *)uchar2_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper uchar2_e_acc_ct0(uchar2_e, cgh); + // CHECK-NEXT: dpct::access_wrapper uchar2_cast_acc_ct1((sycl::uchar2 *)uchar2_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -2221,8 +2221,8 @@ int main_uchar3() { int *uchar3_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper uchar3_e_acc_ct0(uchar3_e, cgh); - // CHECK-NEXT: dpct::access_wrapper uchar3_cast_acc_ct1((sycl::uchar3 *)uchar3_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper uchar3_e_acc_ct0(uchar3_e, cgh); + // CHECK-NEXT: dpct::access_wrapper uchar3_cast_acc_ct1((sycl::uchar3 *)uchar3_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -2293,8 +2293,8 @@ int main_uchar4() { int *uchar4_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper uchar4_e_acc_ct0(uchar4_e, cgh); - // CHECK-NEXT: dpct::access_wrapper uchar4_cast_acc_ct1((sycl::uchar4 *)uchar4_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper uchar4_e_acc_ct0(uchar4_e, cgh); + // CHECK-NEXT: dpct::access_wrapper uchar4_cast_acc_ct1((sycl::uchar4 *)uchar4_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -2365,8 +2365,8 @@ int main_uint1() { int *uint1_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper uint1_e_acc_ct0(uint1_e, cgh); - // CHECK-NEXT: dpct::access_wrapper uint1_cast_acc_ct1((uint32_t *)uint1_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper uint1_e_acc_ct0(uint1_e, cgh); + // CHECK-NEXT: dpct::access_wrapper uint1_cast_acc_ct1((uint32_t *)uint1_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -2437,8 +2437,8 @@ int main_uint2() { int *uint2_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper uint2_e_acc_ct0(uint2_e, cgh); - // CHECK-NEXT: dpct::access_wrapper uint2_cast_acc_ct1((sycl::uint2 *)uint2_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper uint2_e_acc_ct0(uint2_e, cgh); + // CHECK-NEXT: dpct::access_wrapper uint2_cast_acc_ct1((sycl::uint2 *)uint2_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -2509,8 +2509,8 @@ int main_uint3() { int *uint3_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper uint3_e_acc_ct0(uint3_e, cgh); - // CHECK-NEXT: dpct::access_wrapper uint3_cast_acc_ct1((sycl::uint3 *)uint3_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper uint3_e_acc_ct0(uint3_e, cgh); + // CHECK-NEXT: dpct::access_wrapper uint3_cast_acc_ct1((sycl::uint3 *)uint3_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -2581,8 +2581,8 @@ int main_uint4() { int *uint4_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper uint4_e_acc_ct0(uint4_e, cgh); - // CHECK-NEXT: dpct::access_wrapper uint4_cast_acc_ct1((sycl::uint4 *)uint4_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper uint4_e_acc_ct0(uint4_e, cgh); + // CHECK-NEXT: dpct::access_wrapper uint4_cast_acc_ct1((sycl::uint4 *)uint4_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -2653,8 +2653,8 @@ int main_ulong1() { int *ulong1_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper ulong1_e_acc_ct0(ulong1_e, cgh); - // CHECK-NEXT: dpct::access_wrapper ulong1_cast_acc_ct1((uint64_t *)ulong1_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper ulong1_e_acc_ct0(ulong1_e, cgh); + // CHECK-NEXT: dpct::access_wrapper ulong1_cast_acc_ct1((uint64_t *)ulong1_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -2725,8 +2725,8 @@ int main_ulong2() { int *ulong2_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper ulong2_e_acc_ct0(ulong2_e, cgh); - // CHECK-NEXT: dpct::access_wrapper ulong2_cast_acc_ct1((sycl::ulong2 *)ulong2_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper ulong2_e_acc_ct0(ulong2_e, cgh); + // CHECK-NEXT: dpct::access_wrapper ulong2_cast_acc_ct1((sycl::ulong2 *)ulong2_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -2797,8 +2797,8 @@ int main_ulong3() { int *ulong3_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper ulong3_e_acc_ct0(ulong3_e, cgh); - // CHECK-NEXT: dpct::access_wrapper ulong3_cast_acc_ct1((sycl::ulong3 *)ulong3_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper ulong3_e_acc_ct0(ulong3_e, cgh); + // CHECK-NEXT: dpct::access_wrapper ulong3_cast_acc_ct1((sycl::ulong3 *)ulong3_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -2869,8 +2869,8 @@ int main_ulong4() { int *ulong4_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper ulong4_e_acc_ct0(ulong4_e, cgh); - // CHECK-NEXT: dpct::access_wrapper ulong4_cast_acc_ct1((sycl::ulong4 *)ulong4_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper ulong4_e_acc_ct0(ulong4_e, cgh); + // CHECK-NEXT: dpct::access_wrapper ulong4_cast_acc_ct1((sycl::ulong4 *)ulong4_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -2941,8 +2941,8 @@ int main_ulonglong1() { int *ulonglong1_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper ulonglong1_e_acc_ct0(ulonglong1_e, cgh); - // CHECK-NEXT: dpct::access_wrapper ulonglong1_cast_acc_ct1((uint64_t *)ulonglong1_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper ulonglong1_e_acc_ct0(ulonglong1_e, cgh); + // CHECK-NEXT: dpct::access_wrapper ulonglong1_cast_acc_ct1((uint64_t *)ulonglong1_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -3013,8 +3013,8 @@ int main_ulonglong2() { int *ulonglong2_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper ulonglong2_e_acc_ct0(ulonglong2_e, cgh); - // CHECK-NEXT: dpct::access_wrapper ulonglong2_cast_acc_ct1((sycl::ulong2 *)ulonglong2_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper ulonglong2_e_acc_ct0(ulonglong2_e, cgh); + // CHECK-NEXT: dpct::access_wrapper ulonglong2_cast_acc_ct1((sycl::ulong2 *)ulonglong2_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -3085,8 +3085,8 @@ int main_ulonglong3() { int *ulonglong3_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper ulonglong3_e_acc_ct0(ulonglong3_e, cgh); - // CHECK-NEXT: dpct::access_wrapper ulonglong3_cast_acc_ct1((sycl::ulong3 *)ulonglong3_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper ulonglong3_e_acc_ct0(ulonglong3_e, cgh); + // CHECK-NEXT: dpct::access_wrapper ulonglong3_cast_acc_ct1((sycl::ulong3 *)ulonglong3_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -3157,8 +3157,8 @@ int main_ulonglong4() { int *ulonglong4_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper ulonglong4_e_acc_ct0(ulonglong4_e, cgh); - // CHECK-NEXT: dpct::access_wrapper ulonglong4_cast_acc_ct1((sycl::ulong4 *)ulonglong4_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper ulonglong4_e_acc_ct0(ulonglong4_e, cgh); + // CHECK-NEXT: dpct::access_wrapper ulonglong4_cast_acc_ct1((sycl::ulong4 *)ulonglong4_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -3229,8 +3229,8 @@ int main_ushort1() { int *ushort1_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper ushort1_e_acc_ct0(ushort1_e, cgh); - // CHECK-NEXT: dpct::access_wrapper ushort1_cast_acc_ct1((uint16_t *)ushort1_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper ushort1_e_acc_ct0(ushort1_e, cgh); + // CHECK-NEXT: dpct::access_wrapper ushort1_cast_acc_ct1((uint16_t *)ushort1_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -3301,8 +3301,8 @@ int main_ushort2() { int *ushort2_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper ushort2_e_acc_ct0(ushort2_e, cgh); - // CHECK-NEXT: dpct::access_wrapper ushort2_cast_acc_ct1((sycl::ushort2 *)ushort2_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper ushort2_e_acc_ct0(ushort2_e, cgh); + // CHECK-NEXT: dpct::access_wrapper ushort2_cast_acc_ct1((sycl::ushort2 *)ushort2_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -3373,8 +3373,8 @@ int main_ushort3() { int *ushort3_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper ushort3_e_acc_ct0(ushort3_e, cgh); - // CHECK-NEXT: dpct::access_wrapper ushort3_cast_acc_ct1((sycl::ushort3 *)ushort3_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper ushort3_e_acc_ct0(ushort3_e, cgh); + // CHECK-NEXT: dpct::access_wrapper ushort3_cast_acc_ct1((sycl::ushort3 *)ushort3_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -3445,8 +3445,8 @@ int main_ushort4() { int *ushort4_cast; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper ushort4_e_acc_ct0(ushort4_e, cgh); - // CHECK-NEXT: dpct::access_wrapper ushort4_cast_acc_ct1((sycl::ushort4 *)ushort4_cast, cgh); + // CHECK-NEXT: dpct::access_wrapper ushort4_e_acc_ct0(ushort4_e, cgh); + // CHECK-NEXT: dpct::access_wrapper ushort4_cast_acc_ct1((sycl::ushort4 *)ushort4_cast, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), diff --git a/clang/test/dpct/warp_api_outside_inroot/src/kernel_warp.cu b/clang/test/dpct/warp_api_outside_inroot/src/kernel_warp.cu index e9877b2bf56a..09001eb32f69 100644 --- a/clang/test/dpct/warp_api_outside_inroot/src/kernel_warp.cu +++ b/clang/test/dpct/warp_api_outside_inroot/src/kernel_warp.cu @@ -31,7 +31,7 @@ void foo() { //CHECK:dpct::get_out_of_order_queue().submit( //CHECK-NEXT: [&](sycl::handler &cgh) { //CHECK-NEXT: sycl::local_accessor smem_acc_ct1(sycl::range<1>(128), cgh); - //CHECK-NEXT: dpct::access_wrapper input_acc_ct0(input, cgh); + //CHECK-NEXT: dpct::access_wrapper input_acc_ct0(input, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 128), sycl::range<3>(1, 1, 128)), @@ -70,7 +70,7 @@ __global__ void compute_mode(float *input) { void foo_2(float *ptr) { //CHECK:dpct::get_out_of_order_queue().submit( //CHECK-NEXT: [&](sycl::handler &cgh) { - //CHECK-NEXT: dpct::access_wrapper ptr_acc_ct0(ptr, cgh); + //CHECK-NEXT: dpct::access_wrapper ptr_acc_ct0(ptr, cgh); //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)), diff --git a/clang/test/dpct/wildcard_test/abc.cu b/clang/test/dpct/wildcard_test/abc.cu index a1785b2a8586..b5a8676c767c 100644 --- a/clang/test/dpct/wildcard_test/abc.cu +++ b/clang/test/dpct/wildcard_test/abc.cu @@ -19,8 +19,8 @@ int main() { int karg3 = 80; // CHECK: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { - // CHECK-NEXT: dpct::access_wrapper karg2_acc_ct0((const int *)karg2, cgh); - // CHECK-NEXT: dpct::access_wrapper karg2_acc_ct1(karg2, cgh); + // CHECK-NEXT: dpct::access_wrapper karg2_acc_ct0((const int *)karg2, cgh); + // CHECK-NEXT: dpct::access_wrapper karg2_acc_ct1(karg2, cgh); // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(griddim * threaddim, threaddim), From db32444291530e69ca5f142fbce0aa3a9dc5fe5e Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Thu, 10 Oct 2024 07:55:38 +0800 Subject: [PATCH 05/11] Fix helper Signed-off-by: Jiang, Zhiwei --- clang/runtime/dpct-rt/include/dpct/memory.hpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/clang/runtime/dpct-rt/include/dpct/memory.hpp b/clang/runtime/dpct-rt/include/dpct/memory.hpp index 11328725f899..9517fadbff58 100644 --- a/clang/runtime/dpct-rt/include/dpct/memory.hpp +++ b/clang/runtime/dpct-rt/include/dpct/memory.hpp @@ -918,7 +918,7 @@ static buffer_t get_buffer(const void *ptr) { template class access_wrapper { - static_assert(!std::is_same_v, void>, + static_assert(!std::is_same_v, void>, "dataT cannot be void"); sycl::accessor accessor; size_t offset; @@ -954,7 +954,8 @@ static auto get_access(const T *ptr, sycl::handler &cgh) { if (std::is_same_v, void>) return alloc.buffer.get_access(cgh); else - return alloc.buffer.reinterpret(sycl::range<1>(alloc.size / sizeof(T))) + return alloc.buffer + .template reinterpret(sycl::range<1>(alloc.size / sizeof(T))) .get_access(cgh); } else { throw std::runtime_error( From 29f74aea22129370def6f6f1919d079f195b6785 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Thu, 10 Oct 2024 09:17:11 +0800 Subject: [PATCH 06/11] Fix Signed-off-by: Jiang, Zhiwei --- clang/runtime/dpct-rt/include/dpct/memory.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/runtime/dpct-rt/include/dpct/memory.hpp b/clang/runtime/dpct-rt/include/dpct/memory.hpp index 9517fadbff58..916e5f5e413f 100644 --- a/clang/runtime/dpct-rt/include/dpct/memory.hpp +++ b/clang/runtime/dpct-rt/include/dpct/memory.hpp @@ -951,8 +951,8 @@ template , void>) - return alloc.buffer.get_access(cgh); + if constexpr (std::is_same_v, void>) + return alloc.buffer.template get_access(cgh); else return alloc.buffer .template reinterpret(sycl::range<1>(alloc.size / sizeof(T))) From 6b9c98854258926c9ad8d8dbc0256cb853c96980 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Thu, 10 Oct 2024 10:19:48 +0800 Subject: [PATCH 07/11] Fix Signed-off-by: Jiang, Zhiwei --- clang/runtime/dpct-rt/include/dpct/memory.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/runtime/dpct-rt/include/dpct/memory.hpp b/clang/runtime/dpct-rt/include/dpct/memory.hpp index 916e5f5e413f..73076a338eb0 100644 --- a/clang/runtime/dpct-rt/include/dpct/memory.hpp +++ b/clang/runtime/dpct-rt/include/dpct/memory.hpp @@ -956,7 +956,7 @@ static auto get_access(const T *ptr, sycl::handler &cgh) { else return alloc.buffer .template reinterpret(sycl::range<1>(alloc.size / sizeof(T))) - .get_access(cgh); + .template get_access(cgh); } else { throw std::runtime_error( "NULL pointer argument in get_access function is invalid"); From 3406856dded1330a53dfab73a8cb50358ce50d24 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Thu, 10 Oct 2024 13:11:08 +0800 Subject: [PATCH 08/11] Fix Signed-off-by: Jiang, Zhiwei --- clang/runtime/dpct-rt/include/dpct/memory.hpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/clang/runtime/dpct-rt/include/dpct/memory.hpp b/clang/runtime/dpct-rt/include/dpct/memory.hpp index 73076a338eb0..eee2f11d6f0e 100644 --- a/clang/runtime/dpct-rt/include/dpct/memory.hpp +++ b/clang/runtime/dpct-rt/include/dpct/memory.hpp @@ -929,7 +929,8 @@ class access_wrapper { /// \param ptr Pointer to memory. /// \param cgh The command group handler. access_wrapper(const dataT *ptr, sycl::handler &cgh) - : accessor(get_buffer(ptr).get_access(cgh)), offset(0) { + : accessor(get_buffer(ptr).template get_access(cgh)), + offset(0) { auto alloc = detail::mem_mgr::instance().translate_ptr(ptr); offset = (byte_t *)ptr - alloc.alloc_ptr; } @@ -937,7 +938,7 @@ class access_wrapper { /// Get the device pointer. /// /// \returns a device pointer with offset. - dataT get_raw_pointer() const { return (dataT)(&accessor[0] + offset); } + dataT *get_raw_pointer() const { return (dataT *)(&accessor[0] + offset); } }; /// Get the accessor for memory pointed by \p ptr. From 91ecbe4a07d313bfca81777bf27a14f6ad7bb2c2 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Thu, 10 Oct 2024 17:00:30 +0800 Subject: [PATCH 09/11] Update Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/AnalysisInfo.cpp | 19 ++++++++++------ clang/lib/DPCT/AnalysisInfo.h | 3 +-- .../dpct-rt/include/dpct/blas_gemm_utils.hpp | 22 +++++++++---------- clang/runtime/dpct-rt/include/dpct/memory.hpp | 8 +++---- .../dpct-rt/include/dpct/rng_utils.hpp | 4 ++++ 5 files changed, 31 insertions(+), 25 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 23fc808c56d6..03fd859e0cd4 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -5224,10 +5224,11 @@ void DeviceFunctionInfo::mergeTextureObjectList( KernelCallExpr::ArgInfo::ArgInfo(const ParmVarDecl *PVD, KernelArgumentAnalysis &Analysis, const Expr *Arg, bool Used, int Index, - KernelCallExpr *BASE, const ParmVarDecl *TPVD) + KernelCallExpr *BASE) : IsPointer(false), IsRedeclareRequired(false), IsUsedAsLvalueAfterMalloc(Used), Index(Index) { - if (TPVD && TPVD->getType()->isDependentType()) + if (PVD && + PVD->getType()->getTypeClass() == Type::TypeClass::SubstTemplateTypeParm) IsDependentType = true; if (isa(Arg)) { HasImplicitConversion = true; @@ -5868,6 +5869,8 @@ void KernelCallExpr::buildArgsInfo(const CallExpr *CE) { Analysis.setCallSpelling(KCallSpellingRange.first, KCallSpellingRange.second); auto &TexList = getTextureObjectList(); + const auto *FD = CE->getDirectCallee(); + const auto *FTD = FD ? FD->getPrimaryTemplate() : nullptr; for (unsigned Idx = 0; Idx < CE->getNumArgs(); ++Idx) { if (auto Obj = TexList[Idx]) { ArgsInfo.emplace_back(Obj, this); @@ -5876,11 +5879,13 @@ void KernelCallExpr::buildArgsInfo(const CallExpr *CE) { bool Used = true; if (auto *ArgDRE = dyn_cast(Arg->IgnoreImpCasts())) Used = isArgUsedAsLvalueUntil(ArgDRE, CE); - const auto FD = CE->getDirectCallee(); - const FunctionTemplateDecl *FTD = FD ? FD->getPrimaryTemplate() : nullptr; - ArgsInfo.emplace_back( - FD ? FD->parameters()[Idx] : nullptr, Analysis, Arg, Used, Idx, this, - FTD ? FTD->getTemplatedDecl()->parameters()[Idx] : nullptr); + ArgsInfo.emplace_back(FD ? FD->parameters()[Idx] : nullptr, Analysis, Arg, + Used, Idx, this); + if (FTD && FTD->getTemplatedDecl() + ->parameters()[Idx] + ->getType() + ->isDependentType()) + ArgsInfo.back().IsDependentType = true; } } } diff --git a/clang/lib/DPCT/AnalysisInfo.h b/clang/lib/DPCT/AnalysisInfo.h index 285229e6af1a..123d9e3901d3 100644 --- a/clang/lib/DPCT/AnalysisInfo.h +++ b/clang/lib/DPCT/AnalysisInfo.h @@ -2739,8 +2739,7 @@ class KernelCallExpr : public CallFunctionExpr { private: struct ArgInfo { ArgInfo(const ParmVarDecl *PVD, KernelArgumentAnalysis &Analysis, - const Expr *Arg, bool Used, int Index, KernelCallExpr *BASE, - const ParmVarDecl *TPVD); + const Expr *Arg, bool Used, int Index, KernelCallExpr *BASE); ArgInfo(const ParmVarDecl *PVD, const std::string &ArgsArrayName, KernelCallExpr *Kernel); ArgInfo(const ParmVarDecl *PVD, KernelCallExpr *Kernel); diff --git a/clang/runtime/dpct-rt/include/dpct/blas_gemm_utils.hpp b/clang/runtime/dpct-rt/include/dpct/blas_gemm_utils.hpp index 22cfda9dae8a..15e85e80bc7a 100644 --- a/clang/runtime/dpct-rt/include/dpct/blas_gemm_utils.hpp +++ b/clang/runtime/dpct-rt/include/dpct/blas_gemm_utils.hpp @@ -389,8 +389,8 @@ template struct matrix_transform_impl { return q_ptr->submit([&](sycl::handler &cgh) { cgh.depends_on(deps); #ifdef DPCT_USM_LEVEL_NONE - access_wrapper a_acc(a, cgh); - access_wrapper c_acc(c, cgh); + access_wrapper a_acc(static_cast(a), cgh); + access_wrapper c_acc(static_cast(c), cgh); #endif cgh.parallel_for< ::dpct::cs::kernel_name>( @@ -431,7 +431,7 @@ inline sycl::event int2float(::dpct::cs::queue_ptr q_ptr, void *int_ptr, }); } else { return q_ptr->submit([&](sycl::handler &cgh) { - access_wrapper int_acc(int_ptr, cgh); + access_wrapper int_acc(static_cast(int_ptr), cgh); sycl::accessor float_acc(float_buffer, cgh, sycl::write_only, sycl::no_init); cgh.single_task<::dpct::cs::kernel_name>([=]() { @@ -450,8 +450,8 @@ inline sycl::event multiply_impl(::dpct::cs::queue_ptr q_ptr, return q_ptr->submit([&](sycl::handler &cgh) { cgh.depends_on(deps); sycl::accessor result_acc(result, cgh); - access_wrapper a_acc(a, cgh); - access_wrapper b_acc(b, cgh); + access_wrapper a_acc(static_cast(a), cgh); + access_wrapper b_acc(static_cast(b), cgh); cgh.single_task<::dpct::cs::kernel_name>([=]() { auto a_ptr = a_acc.get_raw_pointer(); auto b_ptr = b_acc.get_raw_pointer(); @@ -462,7 +462,7 @@ inline sycl::event multiply_impl(::dpct::cs::queue_ptr q_ptr, return q_ptr->submit([&](sycl::handler &cgh) { cgh.depends_on(deps); sycl::accessor result_acc(result, cgh); - access_wrapper a_acc(a, cgh); + access_wrapper a_acc(static_cast(a), cgh); cgh.single_task<::dpct::cs::kernel_name>([=]() { auto a_ptr = a_acc.get_raw_pointer(); result_acc[0] = result_acc[0] * a_ptr[0]; @@ -472,7 +472,7 @@ inline sycl::event multiply_impl(::dpct::cs::queue_ptr q_ptr, return q_ptr->submit([&](sycl::handler &cgh) { cgh.depends_on(deps); sycl::accessor result_acc(result, cgh); - access_wrapper b_acc(b, cgh); + access_wrapper b_acc(static_cast(b), cgh); cgh.single_task<::dpct::cs::kernel_name>([=]() { auto b_ptr = b_acc.get_raw_pointer(); result_acc[0] = result_acc[0] * b_ptr[0]; @@ -492,7 +492,7 @@ template struct scale_d_impl { cgh.depends_on(deps); access_wrapper d_scale_acc( static_cast(d_scale_ptr), cgh); - access_wrapper d_acc(d, cgh); + access_wrapper d_acc(static_cast(d), cgh); cgh.parallel_for<::dpct::cs::kernel_name>( sycl::range<2>(ld, cols), [=](sycl::id<2> idx) { float scale_factor = d_scale_acc.get_raw_pointer()[0]; @@ -511,7 +511,7 @@ template struct scale_d_impl { cgh.depends_on(deps); access_wrapper d_scale_acc( static_cast(d_scale_ptr), cgh); - access_wrapper d_acc(d, cgh); + access_wrapper d_acc(static_cast(d), cgh); cgh.parallel_for<::dpct::cs::kernel_name>( sycl::range<2>(ld, cols), [=](sycl::id<2> idx) { float scale_factor = @@ -640,7 +640,7 @@ template struct absmax_impl { auto absmax_reduction = sycl::reduction( get_buffer(absmax_ptr), cgh, T(0), abs_max_op(), {sycl::property::reduction::initialize_to_identity()}); - access_wrapper new_d_acc(new_d, cgh); + access_wrapper new_d_acc(static_cast(new_d), cgh); #else auto absmax_reduction = sycl::reduction( (T *)(absmax_ptr), T(0), abs_max_op(), @@ -969,7 +969,7 @@ inline sycl::event matmul(descriptor_ptr handle, matmul_desc_ptr compute_desc, auto buf = ::dnnl::sycl_interop::get_buffer(*scales_alpha); if (dpct::is_device_ptr(alpha)) { scalar_alpha_e = q_ptr->submit([&](sycl::handler &cgh) { - access_wrapper alpha_acc(alpha, cgh); + access_wrapper alpha_acc(static_cast(alpha), cgh); sycl::accessor acc(buf, cgh, sycl::write_only, sycl::no_init); cgh.single_task<::dpct::cs::kernel_name>( [=]() { acc[0] = alpha_acc.get_raw_pointer()[0]; }); diff --git a/clang/runtime/dpct-rt/include/dpct/memory.hpp b/clang/runtime/dpct-rt/include/dpct/memory.hpp index eee2f11d6f0e..72576aa94d7d 100644 --- a/clang/runtime/dpct-rt/include/dpct/memory.hpp +++ b/clang/runtime/dpct-rt/include/dpct/memory.hpp @@ -915,11 +915,9 @@ static buffer_t get_buffer(const void *ptr) { } /// A wrapper class contains an accessor and an offset. -template class access_wrapper { - static_assert(!std::is_same_v, void>, - "dataT cannot be void"); sycl::accessor accessor; size_t offset; @@ -928,7 +926,7 @@ class access_wrapper { /// /// \param ptr Pointer to memory. /// \param cgh The command group handler. - access_wrapper(const dataT *ptr, sycl::handler &cgh) + access_wrapper(PtrT ptr, sycl::handler &cgh) : accessor(get_buffer(ptr).template get_access(cgh)), offset(0) { auto alloc = detail::mem_mgr::instance().translate_ptr(ptr); @@ -938,7 +936,7 @@ class access_wrapper { /// Get the device pointer. /// /// \returns a device pointer with offset. - dataT *get_raw_pointer() const { return (dataT *)(&accessor[0] + offset); } + PtrT get_raw_pointer() const { return (PtrT)(&accessor[0] + offset); } }; /// Get the accessor for memory pointed by \p ptr. diff --git a/clang/runtime/dpct-rt/include/dpct/rng_utils.hpp b/clang/runtime/dpct-rt/include/dpct/rng_utils.hpp index 4a2460763b5a..d741739cfdaf 100644 --- a/clang/runtime/dpct-rt/include/dpct/rng_utils.hpp +++ b/clang/runtime/dpct-rt/include/dpct/rng_utils.hpp @@ -614,4 +614,8 @@ create_host_rng(const random_engine_type type, } // namespace rng } // namespace dpct +template +struct sycl::is_device_copyable> + : std::true_type {}; + #endif // __DPCT_RNG_UTILS_HPP__ From c4a47b8c665df086892d371ba2f6813b0801deeb Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Fri, 11 Oct 2024 09:03:12 +0800 Subject: [PATCH 10/11] Fix lit Signed-off-by: Jiang, Zhiwei --- clang/test/dpct/texture_layered.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/dpct/texture_layered.cu b/clang/test/dpct/texture_layered.cu index cfd91639d2ac..30d928175b4e 100644 --- a/clang/test/dpct/texture_layered.cu +++ b/clang/test/dpct/texture_layered.cu @@ -116,7 +116,7 @@ int main() { // CHECK-NEXT: cgh.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: kernel((sycl::float4 *)(&d_acc_ct0[0]), dpct::image_accessor_ext(tex42_smpl, tex42_acc), dpct::image_accessor_ext(tex21_smpl, tex21_acc)); + // CHECK-NEXT: kernel(&d_acc_ct0[0], dpct::image_accessor_ext(tex42_smpl, tex42_acc), dpct::image_accessor_ext(tex21_smpl, tex21_acc)); // CHECK-NEXT: }); // CHECK-NEXT: }); kernel<<<1, 1>>>(d); From ea5ed53867bc1788926758e9efd7ae8abc84ac4d Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Mon, 14 Oct 2024 08:46:42 +0800 Subject: [PATCH 11/11] Update Signed-off-by: Jiang, Zhiwei --- .../dpct-rt/include/dpct/blas_gemm_utils.hpp | 22 +++++++++---------- clang/runtime/dpct-rt/include/dpct/memory.hpp | 11 +++++++--- 2 files changed, 19 insertions(+), 14 deletions(-) diff --git a/clang/runtime/dpct-rt/include/dpct/blas_gemm_utils.hpp b/clang/runtime/dpct-rt/include/dpct/blas_gemm_utils.hpp index 15e85e80bc7a..22cfda9dae8a 100644 --- a/clang/runtime/dpct-rt/include/dpct/blas_gemm_utils.hpp +++ b/clang/runtime/dpct-rt/include/dpct/blas_gemm_utils.hpp @@ -389,8 +389,8 @@ template struct matrix_transform_impl { return q_ptr->submit([&](sycl::handler &cgh) { cgh.depends_on(deps); #ifdef DPCT_USM_LEVEL_NONE - access_wrapper a_acc(static_cast(a), cgh); - access_wrapper c_acc(static_cast(c), cgh); + access_wrapper a_acc(a, cgh); + access_wrapper c_acc(c, cgh); #endif cgh.parallel_for< ::dpct::cs::kernel_name>( @@ -431,7 +431,7 @@ inline sycl::event int2float(::dpct::cs::queue_ptr q_ptr, void *int_ptr, }); } else { return q_ptr->submit([&](sycl::handler &cgh) { - access_wrapper int_acc(static_cast(int_ptr), cgh); + access_wrapper int_acc(int_ptr, cgh); sycl::accessor float_acc(float_buffer, cgh, sycl::write_only, sycl::no_init); cgh.single_task<::dpct::cs::kernel_name>([=]() { @@ -450,8 +450,8 @@ inline sycl::event multiply_impl(::dpct::cs::queue_ptr q_ptr, return q_ptr->submit([&](sycl::handler &cgh) { cgh.depends_on(deps); sycl::accessor result_acc(result, cgh); - access_wrapper a_acc(static_cast(a), cgh); - access_wrapper b_acc(static_cast(b), cgh); + access_wrapper a_acc(a, cgh); + access_wrapper b_acc(b, cgh); cgh.single_task<::dpct::cs::kernel_name>([=]() { auto a_ptr = a_acc.get_raw_pointer(); auto b_ptr = b_acc.get_raw_pointer(); @@ -462,7 +462,7 @@ inline sycl::event multiply_impl(::dpct::cs::queue_ptr q_ptr, return q_ptr->submit([&](sycl::handler &cgh) { cgh.depends_on(deps); sycl::accessor result_acc(result, cgh); - access_wrapper a_acc(static_cast(a), cgh); + access_wrapper a_acc(a, cgh); cgh.single_task<::dpct::cs::kernel_name>([=]() { auto a_ptr = a_acc.get_raw_pointer(); result_acc[0] = result_acc[0] * a_ptr[0]; @@ -472,7 +472,7 @@ inline sycl::event multiply_impl(::dpct::cs::queue_ptr q_ptr, return q_ptr->submit([&](sycl::handler &cgh) { cgh.depends_on(deps); sycl::accessor result_acc(result, cgh); - access_wrapper b_acc(static_cast(b), cgh); + access_wrapper b_acc(b, cgh); cgh.single_task<::dpct::cs::kernel_name>([=]() { auto b_ptr = b_acc.get_raw_pointer(); result_acc[0] = result_acc[0] * b_ptr[0]; @@ -492,7 +492,7 @@ template struct scale_d_impl { cgh.depends_on(deps); access_wrapper d_scale_acc( static_cast(d_scale_ptr), cgh); - access_wrapper d_acc(static_cast(d), cgh); + access_wrapper d_acc(d, cgh); cgh.parallel_for<::dpct::cs::kernel_name>( sycl::range<2>(ld, cols), [=](sycl::id<2> idx) { float scale_factor = d_scale_acc.get_raw_pointer()[0]; @@ -511,7 +511,7 @@ template struct scale_d_impl { cgh.depends_on(deps); access_wrapper d_scale_acc( static_cast(d_scale_ptr), cgh); - access_wrapper d_acc(static_cast(d), cgh); + access_wrapper d_acc(d, cgh); cgh.parallel_for<::dpct::cs::kernel_name>( sycl::range<2>(ld, cols), [=](sycl::id<2> idx) { float scale_factor = @@ -640,7 +640,7 @@ template struct absmax_impl { auto absmax_reduction = sycl::reduction( get_buffer(absmax_ptr), cgh, T(0), abs_max_op(), {sycl::property::reduction::initialize_to_identity()}); - access_wrapper new_d_acc(static_cast(new_d), cgh); + access_wrapper new_d_acc(new_d, cgh); #else auto absmax_reduction = sycl::reduction( (T *)(absmax_ptr), T(0), abs_max_op(), @@ -969,7 +969,7 @@ inline sycl::event matmul(descriptor_ptr handle, matmul_desc_ptr compute_desc, auto buf = ::dnnl::sycl_interop::get_buffer(*scales_alpha); if (dpct::is_device_ptr(alpha)) { scalar_alpha_e = q_ptr->submit([&](sycl::handler &cgh) { - access_wrapper alpha_acc(static_cast(alpha), cgh); + access_wrapper alpha_acc(alpha, cgh); sycl::accessor acc(buf, cgh, sycl::write_only, sycl::no_init); cgh.single_task<::dpct::cs::kernel_name>( [=]() { acc[0] = alpha_acc.get_raw_pointer()[0]; }); diff --git a/clang/runtime/dpct-rt/include/dpct/memory.hpp b/clang/runtime/dpct-rt/include/dpct/memory.hpp index 72576aa94d7d..be5950d61d72 100644 --- a/clang/runtime/dpct-rt/include/dpct/memory.hpp +++ b/clang/runtime/dpct-rt/include/dpct/memory.hpp @@ -926,12 +926,17 @@ class access_wrapper { /// /// \param ptr Pointer to memory. /// \param cgh The command group handler. - access_wrapper(PtrT ptr, sycl::handler &cgh) - : accessor(get_buffer(ptr).template get_access(cgh)), - offset(0) { + access_wrapper(const void *ptr, sycl::handler &cgh) + : accessor(get_buffer(ptr).get_access(cgh)), offset(0) { auto alloc = detail::mem_mgr::instance().translate_ptr(ptr); offset = (byte_t *)ptr - alloc.alloc_ptr; } + template + access_wrapper( + PtrT ptr, sycl::handler &cgh, + typename std::enable_if_t>, void *>> * = 0) + : access_wrapper((const void *)ptr, cgh) {} /// Get the device pointer. ///