Skip to content

Commit 31d0bbf

Browse files
zhiweij1zhiweij1
authored
[SYCLomatic] Add "template" keyword before method for template object (#1948)
 Signed-off-by: Jiang, Zhiwei <[email protected]>
1 parent 0566d11 commit 31d0bbf

13 files changed

+66
-52
lines changed

‎clang/lib/DPCT/AnalysisInfo.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3088,7 +3088,10 @@ ParameterStream &MemVarInfo::getKernelArg(ParameterStream &PS) {
30883088
if (AccMode == Pointer) {
30893089
if (!getType()->isWritten())
30903090
PS << "(" << getAccessorDataType(false, true) << " *)";
3091-
PS << getAccessorName() << ".get_multi_ptr<" << MapNames::getClNamespace()
3091+
PS << getAccessorName() << ".";
3092+
if (getType()->isTemplate())
3093+
PS << "template ";
3094+
PS << "get_multi_ptr<" << MapNames::getClNamespace()
30923095
<< "access::decorated::no>().get()";
30933096
} else {
30943097
PS << getAccessorName();

‎clang/test/dpct/kernel-call.cu

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -682,7 +682,7 @@ void test_ctor() {
682682
//CHECK-NEXT: cgh.parallel_for<dpct_kernel_name<class k11_{{[0-9a-z]+}}, TT>>(
683683
//CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
684684
//CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) {
685-
//CHECK-NEXT: k11<TT>(a, temp_ct1_acc_ct1.get_multi_ptr<sycl::access::decorated::no>().get(), temp2_ct1_acc_ct1.get_multi_ptr<sycl::access::decorated::no>().get());
685+
//CHECK-NEXT: k11<TT>(a, temp_ct1_acc_ct1.template get_multi_ptr<sycl::access::decorated::no>().get(), temp2_ct1_acc_ct1.template get_multi_ptr<sycl::access::decorated::no>().get());
686686
//CHECK-NEXT: });
687687
//CHECK-NEXT: });
688688
//CHECK-NEXT:}
@@ -731,7 +731,7 @@ void foo11() {
731731
//CHECK-NEXT: cgh.parallel_for<dpct_kernel_name<class k12_{{[0-9a-z]+}}, TT>>(
732732
//CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
733733
//CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) {
734-
//CHECK-NEXT: k12<TT>(a, temp_ct1_acc_ct1.get_multi_ptr<sycl::access::decorated::no>().get(), temp2_ct1_acc_ct1.get_multi_ptr<sycl::access::decorated::no>().get());
734+
//CHECK-NEXT: k12<TT>(a, temp_ct1_acc_ct1.template get_multi_ptr<sycl::access::decorated::no>().get(), temp2_ct1_acc_ct1.template get_multi_ptr<sycl::access::decorated::no>().get());
735735
//CHECK-NEXT: });
736736
//CHECK-NEXT: });
737737
//CHECK-NEXT:}

‎clang/test/dpct/kernel_without_name.cu

Lines changed: 12 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -325,7 +325,7 @@ void run_foo8() {
325325
// CHECK-NEXT: cgh.parallel_for(
326326
// CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
327327
// CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) {
328-
// CHECK-NEXT: foo_kernel7<T>(i, i, mem_acc_ct1.get_multi_ptr<sycl::access::decorated::no>().get());
328+
// CHECK-NEXT: foo_kernel7<T>(i, i, mem_acc_ct1.template get_multi_ptr<sycl::access::decorated::no>().get());
329329
// CHECK-NEXT: });
330330
// CHECK-NEXT: });
331331
int i;
@@ -358,3 +358,14 @@ void run_foo9() {
358358
int i;
359359
foo_kernel8<<<1, 1>>>(i, i);
360360
}
361+
362+
template <typename T> __global__ void foo_kernel9() { __shared__ T mem[256]; }
363+
364+
template <typename T> void run_foo10() {
365+
// CHECK: cgh.parallel_for(
366+
// CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
367+
// CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) {
368+
// CHECK-NEXT: foo_kernel9<T>(mem_acc_ct1.template get_multi_ptr<sycl::access::decorated::no>().get());
369+
// CHECK-NEXT: });
370+
foo_kernel9<T><<<1, 1>>>();
371+
}

‎clang/test/dpct/launch-kernel-cooperative-usm.cu

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -75,7 +75,7 @@ int main() {
7575
// CHECK-NEXT: cgh.parallel_for(
7676
// CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 16), sycl::range<3>(1, 1, 16)),
7777
// CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) {
78-
// CHECK-NEXT: template_kernel<int>(d_ct0, item_ct1, dpct_local_acc_ct1.get_multi_ptr<sycl::access::decorated::no>().get(), s_acc_ct1.get_multi_ptr<sycl::access::decorated::no>().get());
78+
// CHECK-NEXT: template_kernel<int>(d_ct0, item_ct1, dpct_local_acc_ct1.template get_multi_ptr<sycl::access::decorated::no>().get(), s_acc_ct1.template get_multi_ptr<sycl::access::decorated::no>().get());
7979
// CHECK-NEXT: });
8080
// CHECK-NEXT: });
8181
cudaLaunchCooperativeKernel((const void *)&template_kernel<int>, dim3(16), dim3(16), args, 32, stream);

‎clang/test/dpct/launch-kernel-cooperative.cu

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -74,7 +74,7 @@ int main() {
7474
// CHECK-NEXT: cgh.parallel_for(
7575
// CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 16), sycl::range<3>(1, 1, 16)),
7676
// CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) {
77-
// CHECK-NEXT: template_kernel<int>(d_acc_ct0.get_raw_pointer(), item_ct1, dpct_local_acc_ct1.get_multi_ptr<sycl::access::decorated::no>().get(), s_acc_ct1.get_multi_ptr<sycl::access::decorated::no>().get());
77+
// CHECK-NEXT: template_kernel<int>(d_acc_ct0.get_raw_pointer(), item_ct1, dpct_local_acc_ct1.template get_multi_ptr<sycl::access::decorated::no>().get(), s_acc_ct1.template get_multi_ptr<sycl::access::decorated::no>().get());
7878
// CHECK-NEXT: });
7979
// CHECK-NEXT: });
8080
cudaLaunchCooperativeKernel((const void *)&template_kernel<int>, dim3(16), dim3(16), args, 32, stream);

‎clang/test/dpct/launch-kernel-usm.cu

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -72,7 +72,7 @@ int main() {
7272
// CHECK-NEXT: cgh.parallel_for(
7373
// CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 16), sycl::range<3>(1, 1, 16)),
7474
// CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) {
75-
// CHECK-NEXT: template_kernel<int>(d_ct0, item_ct1, dpct_local_acc_ct1.get_multi_ptr<sycl::access::decorated::no>().get(), s_acc_ct1.get_multi_ptr<sycl::access::decorated::no>().get());
75+
// CHECK-NEXT: template_kernel<int>(d_ct0, item_ct1, dpct_local_acc_ct1.template get_multi_ptr<sycl::access::decorated::no>().get(), s_acc_ct1.template get_multi_ptr<sycl::access::decorated::no>().get());
7676
// CHECK-NEXT: });
7777
// CHECK-NEXT: });
7878
cudaLaunchKernel((const void *)&template_kernel<int>, dim3(16), dim3(16), args, 32, stream);

‎clang/test/dpct/launch-kernel.cu

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -72,7 +72,7 @@ int main() {
7272
// CHECK-NEXT: cgh.parallel_for(
7373
// CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 16), sycl::range<3>(1, 1, 16)),
7474
// CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) {
75-
// CHECK-NEXT: template_kernel<int>(d_acc_ct0.get_raw_pointer(), item_ct1, dpct_local_acc_ct1.get_multi_ptr<sycl::access::decorated::no>().get(), s_acc_ct1.get_multi_ptr<sycl::access::decorated::no>().get());
75+
// CHECK-NEXT: template_kernel<int>(d_acc_ct0.get_raw_pointer(), item_ct1, dpct_local_acc_ct1.template get_multi_ptr<sycl::access::decorated::no>().get(), s_acc_ct1.template get_multi_ptr<sycl::access::decorated::no>().get());
7676
// CHECK-NEXT: });
7777
// CHECK-NEXT: });
7878
cudaLaunchKernel((const void *)&template_kernel<int>, dim3(16), dim3(16), args, 32, stream);

‎clang/test/dpct/sharedmem_var_dynamic.cu

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -48,7 +48,7 @@ void testTemplate(){
4848
// CHECK-NEXT: cgh.parallel_for<dpct_kernel_name<class templateReverse_{{[a-f0-9]+}}, T>>(
4949
// CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, n), sycl::range<3>(1, 1, n)),
5050
// CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) {
51-
// CHECK-NEXT: templateReverse<T>(d_d_acc_ct0.get_raw_pointer(), n, item_ct1, dpct_local_acc_ct1.get_multi_ptr<sycl::access::decorated::no>().get());
51+
// CHECK-NEXT: templateReverse<T>(d_d_acc_ct0.get_raw_pointer(), n, item_ct1, dpct_local_acc_ct1.template get_multi_ptr<sycl::access::decorated::no>().get());
5252
// CHECK-NEXT: });
5353
// CHECK-NEXT: });
5454
templateReverse<T><<<1, n, mem_size>>>(d_d, n);
@@ -101,7 +101,7 @@ int main(void) {
101101
// CHECK-NEXT: cgh.parallel_for<dpct_kernel_name<class templateReverse_{{[a-f0-9]+}}, int>>(
102102
// CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, n), sycl::range<3>(1, 1, n)),
103103
// CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) {
104-
// CHECK-NEXT: templateReverse<int>((int *)(&d_d_acc_ct0[0]), n, item_ct1, dpct_local_acc_ct1.get_multi_ptr<sycl::access::decorated::no>().get());
104+
// CHECK-NEXT: templateReverse<int>((int *)(&d_d_acc_ct0[0]), n, item_ct1, dpct_local_acc_ct1.template get_multi_ptr<sycl::access::decorated::no>().get());
105105
// CHECK-NEXT: });
106106
// CHECK-NEXT: });
107107
templateReverse<int><<<1, n, 4>>>(d_d, n);

0 commit comments

Comments
 (0)