Skip to content

Commit d0d6a24

Browse files
intwanghaointwanghao
authored
[SYCLomatic] Fix two lit cases make sure themselves can be built #815
Signed-off-by: Hao Wang <[email protected]>
1 parent 080b770 commit d0d6a24

File tree

2 files changed

+43
-32
lines changed

2 files changed

+43
-32
lines changed

‎clang/test/dpct/predefined_macro_replacement.cu‎

Lines changed: 27 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -4,53 +4,55 @@
44
//CHECK: #ifdef DPCT_COMPATIBILITY_TEMP
55
//CHECK-NEXT: void hello(const sycl::stream &[[STREAM:stream_ct1]]) { [[STREAM]] << "foo"; }
66
#ifdef __CUDA_ARCH__
7-
__global__ void hello() { printf("foo"); }
7+
__device__ void hello() { printf("foo"); }
88
#else
9-
__global__ void hello() { printf("other"); }
9+
__device__ void hello() { printf("other"); }
1010
#endif
1111

1212
//CHECK: #ifndef DPCT_COMPATIBILITY_TEMP
1313
#ifndef __NVCC__
14-
__global__ void hello2() { printf("hello2"); }
14+
__device__ void hello2() { printf("hello2"); }
1515
#endif
1616
//CHECK: #if defined(SYCL_LANGUAGE_VERSION)
1717
#if defined(__CUDACC__)
18-
__global__ void hello3() { printf("hello2"); }
18+
__device__ void hello3() { printf("hello2"); }
1919
#endif
2020

2121
#if defined(xxx)
22-
__global__ void hello4() { printf("hello2"); }
22+
__device__ void hello4() { printf("hello2"); }
2323
//CHECK: #elif defined(DPCT_COMPATIBILITY_TEMP)
2424
//CHECK-NEXT: void hello4(const sycl::stream &[[STREAM]]) { [[STREAM]] << "hello2"; }
2525
#elif defined(__CUDA_ARCH__)
26-
__global__ void hello4() { printf("hello2"); }
26+
__device__ void hello4() { printf("hello2"); }
2727
#endif
2828

2929
#if defined(xxx)
30-
__global__ void hello5() { printf("hello2"); }
30+
__device__ void hello5() { printf("hello2"); }
3131
//CHECK: #elif (DPCT_COMPATIBILITY_TEMP >= 400)
3232
//CHECK-NEXT: void hello5(const sycl::stream &[[STREAM]]) { [[STREAM]] << "hello2"; }
3333
#elif (__CUDA_ARCH__ >= 400)
34-
__global__ void hello5() { printf("hello2"); }
34+
__device__ void hello5() { printf("hello2"); }
3535
#endif
3636

3737
//CHECK: #if defined(DPCT_COMPATIBILITY_TEMP)
3838
//CHECK-NEXT: void hello6(const sycl::stream &[[STREAM]]) { [[STREAM]] << "hello2"; }
3939
#if defined(__CUDA_ARCH__)
40-
__global__ void hello6() { printf("hello2"); }
40+
__device__ void hello6() { printf("hello2"); }
4141
#endif
4242

4343
//CHECK: #ifndef DPCT_COMPATIBILITY_TEMP
44-
//CHECK-NEXT: __global__ void hello7() { printf("hello2"); }
44+
//CHECK-NEXT: __device__ void hello7() { printf("hello2"); }
4545
//CHECK-NEXT: #else
4646
//CHECK-NEXT: void hello7(const sycl::stream &[[STREAM]]) { [[STREAM]] << "hello2"; }
4747
#ifndef __CUDA_ARCH__
48-
__global__ void hello7() { printf("hello2"); }
48+
__device__ void hello7() { printf("hello2"); }
4949
#else
50-
__global__ void hello7() { printf("hello2"); }
50+
__device__ void hello7() { printf("hello2"); }
5151
#endif
5252

53-
__global__ void test(){
53+
__global__ void hello8() {}
54+
55+
__device__ void test(){
5456
//CHECK:#if (DPCT_COMPATIBILITY_TEMP >= 400) && (DPCT_COMPATIBILITY_TEMP >= 400)
5557
//CHECK-NEXT:[[STREAM]] << ">400, \n";
5658
//CHECK-NEXT:#elif (DPCT_COMPATIBILITY_TEMP >200)
@@ -70,25 +72,25 @@ printf("<200, \n");
7072

7173
int main() {
7274
//CHECK: #if defined(DPCT_COMPATIBILITY_TEMP)
73-
//CHECK-NEXT: q_ct1.submit(
75+
//CHECK-NEXT: q_ct1.parallel_for(
7476
#if defined(__NVCC__)
75-
hello<<<1,1>>>();
77+
hello8<<<1,1>>>();
7678
#else
7779
hello();
7880
#endif
7981

8082
//CHECK: #ifdef DPCT_COMPATIBILITY_TEMP
81-
//CHECK-NEXT: q_ct1.submit(
83+
//CHECK-NEXT: q_ct1.parallel_for(
8284
#ifdef __NVCC__
83-
hello<<<1,1>>>();
85+
hello8<<<1,1>>>();
8486
#else
8587
hello();
8688
#endif
8789

8890
//CHECK: #if DPCT_COMPATIBILITY_TEMP
89-
//CHECK-NEXT: q_ct1.submit(
91+
//CHECK-NEXT: q_ct1.parallel_for(
9092
#if __NVCC__
91-
hello<<<1,1>>>();
93+
hello8<<<1,1>>>();
9294
#else
9395
hello();
9496
#endif
@@ -131,6 +133,9 @@ int foo(int num) {
131133
#else
132134
cudaThreadExit();
133135
#endif
136+
137+
return 0;
138+
134139
}
135140

136141
int foo1() {
@@ -146,4 +151,6 @@ int2 a;
146151
#ifndef CUDART_VERSION
147152
int2 b;
148153
#endif
149-
}
154+
a.x = 1;
155+
return a.x;
156+
}

‎clang/test/dpct/user_defined_cuda_arch.cu‎

Lines changed: 16 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -6,44 +6,44 @@
66
//CHECK: #ifdef DPCT_COMPATIBILITY_TEMP
77
//CHECK-NEXT: void hello(const sycl::stream &[[STREAM:stream_ct1]]) { [[STREAM]] << "foo"; }
88
#ifdef __CUDA_ARCH__
9-
__global__ void hello() { printf("foo"); }
9+
__device__ void hello() { printf("foo"); }
1010
#else
11-
__global__ void hello() { printf("other"); }
11+
__device__ void hello() { printf("other"); }
1212
#endif
1313

1414
#if defined(xxx)
15-
__global__ void hello4() { printf("hello2"); }
15+
__device__ void hello4() { printf("hello2"); }
1616
//CHECK: #elif defined(DPCT_COMPATIBILITY_TEMP)
1717
//CHECK-NEXT: void hello4(const sycl::stream &[[STREAM]]) { [[STREAM]] << "hello2"; }
1818
#elif defined(__CUDA_ARCH__)
19-
__global__ void hello4() { printf("hello2"); }
19+
__device__ void hello4() { printf("hello2"); }
2020
#endif
2121

2222
#if defined(xxx)
23-
__global__ void hello5() { printf("hello2"); }
23+
__device__ void hello5() { printf("hello2"); }
2424
//CHECK: #elif (DPCT_COMPATIBILITY_TEMP >= 400)
25-
//CHECK-NEXT: __global__ void hello5() { printf("hello2"); }
25+
//CHECK-NEXT: __device__ void hello5() { printf("hello2"); }
2626
#elif (__CUDA_ARCH__ >= 400)
27-
__global__ void hello5() { printf("hello2"); }
27+
__device__ void hello5() { printf("hello2"); }
2828
#endif
2929

3030
//CHECK: #if defined(DPCT_COMPATIBILITY_TEMP)
3131
//CHECK-NEXT: void hello6(const sycl::stream &[[STREAM]]) { [[STREAM]] << "hello2"; }
3232
#if defined(__CUDA_ARCH__)
33-
__global__ void hello6() { printf("hello2"); }
33+
__device__ void hello6() { printf("hello2"); }
3434
#endif
3535

3636
//CHECK: #ifndef DPCT_COMPATIBILITY_TEMP
37-
//CHECK-NEXT: __global__ void hello7() { printf("hello2"); }
37+
//CHECK-NEXT: __device__ void hello7() { printf("hello2"); }
3838
//CHECK-NEXT: #else
3939
//CHECK-NEXT: void hello7(const sycl::stream &[[STREAM]]) { [[STREAM]] << "hello2"; }
4040
#ifndef __CUDA_ARCH__
41-
__global__ void hello7() { printf("hello2"); }
41+
__device__ void hello7() { printf("hello2"); }
4242
#else
43-
__global__ void hello7() { printf("hello2"); }
43+
__device__ void hello7() { printf("hello2"); }
4444
#endif
4545

46-
__global__ void test(){
46+
__device__ void test(){
4747
//CHECK:#if (DPCT_COMPATIBILITY_TEMP >= 400) && (DPCT_COMPATIBILITY_TEMP >= 400)
4848
//CHECK-NEXT:printf(">400, \n");
4949
//CHECK-NEXT:#elif (DPCT_COMPATIBILITY_TEMP >200)
@@ -59,3 +59,7 @@ printf(">200, \n");
5959
printf("<200, \n");
6060
#endif
6161
}
62+
63+
int main(){
64+
return 0;
65+
}

0 commit comments

Comments
 (0)