1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
|
// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --hip-kernel-execution-syntax %clang_args
#include <iostream>
#include <algorithm>
// CHECK: #include <hip/hip_runtime.h>
#include <cuda.h>
template<typename T>
__global__ void axpy(T a, T *x, T *y) {
y[threadIdx.x] = a * x[threadIdx.x];
}
template<typename T1, typename T2>
__global__ void axpy_2(T1 a, T2 *x, T2 *y) {
y[threadIdx.x] = a * x[threadIdx.x];
}
template<typename T>
__global__ void axpy_empty() {
}
__global__ void empty() {
}
__global__ void nonempty(int x, int y, int z) {
}
int main(int argc, char* argv[]) {
const int kDataLen = 4;
float a = 2.0f;
float host_x[kDataLen] = {1.0f, 2.0f, 3.0f, 4.0f};
float host_y[kDataLen];
// Copy input data to device.
float* device_x;
float* device_y;
// CHECK: hipMalloc(&device_x, kDataLen * sizeof(float));
cudaMalloc(&device_x, kDataLen * sizeof(float));
// CHECK: hipMalloc(&device_y, kDataLen * sizeof(float));
cudaMalloc(&device_y, kDataLen * sizeof(float));
// CHECK: hipMemcpy(device_x, host_x, kDataLen * sizeof(float), hipMemcpyHostToDevice);
cudaMemcpy(device_x, host_x, kDataLen * sizeof(float), cudaMemcpyHostToDevice);
int x = 1, y = 2, z = 3;
size_t N = 32;
// CHECK: hipStream_t stream = NULL;
cudaStream_t stream = NULL;
// CHECK: hipStreamCreate(&stream);
cudaStreamCreate(&stream);
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy<float>), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y);
axpy<float><<<1, kDataLen>>>(a, device_x, device_y);
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy<float>), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y);
axpy<float><<<dim3(1), kDataLen>>>(a, device_x, device_y);
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy<float>), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y);
axpy<float><<<1, dim3(kDataLen)>>>(a, device_x, device_y);
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy<float>), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y);
axpy<float><<<dim3(1), dim3(kDataLen)>>>(a, device_x, device_y);
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy<float>), dim3(1), dim3(kDataLen), N, 0, a, device_x, device_y);
axpy<float><<<1, kDataLen, N>>>(a, device_x, device_y);
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy<float>), dim3(1), dim3(kDataLen), N, 0, a, device_x, device_y);
axpy<float><<<dim3(1), kDataLen, N>>>(a, device_x, device_y);
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy<float>), dim3(1), dim3(kDataLen), N, 0, a, device_x, device_y);
axpy<float><<<1, dim3(kDataLen), N>>>(a, device_x, device_y);
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy<float>), dim3(1), dim3(kDataLen), N, 0, a, device_x, device_y);
axpy<float><<<dim3(1), dim3(kDataLen), N>>>(a, device_x, device_y);
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy<float>), dim3(1), dim3(kDataLen), N, stream, a, device_x, device_y);
axpy<float><<<1, kDataLen, N, stream>>>(a, device_x, device_y);
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy<float>), dim3(1), dim3(kDataLen), N, stream, a, device_x, device_y);
axpy<float><<<dim3(1), kDataLen, N, stream>>>(a, device_x, device_y);
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy<float>), dim3(1), dim3(kDataLen), N, stream, a, device_x, device_y);
axpy<float><<<1, dim3(kDataLen), N, stream>>>(a, device_x, device_y);
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy<float>), dim3(1), dim3(kDataLen), N, stream, a, device_x, device_y);
axpy<float><<<dim3(1), dim3(kDataLen), N, stream>>>(a, device_x, device_y);
double h_x[kDataLen] = {1.0f, 2.0f, 3.0f, 4.0f};
double h_y[kDataLen];
// Copy input data to device.
double* d_x;
double* d_y;
// CHECK: hipMalloc(&d_x, kDataLen * sizeof(double));
cudaMalloc(&d_x, kDataLen * sizeof(double));
// CHECK: hipMalloc(&d_y, kDataLen * sizeof(double));
cudaMalloc(&d_y, kDataLen * sizeof(double));
// CHECK: hipMemcpy(d_x, h_x, kDataLen * sizeof(double), hipMemcpyHostToDevice);
cudaMemcpy(d_x, h_x, kDataLen * sizeof(double), cudaMemcpyHostToDevice);
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_2<float,double>), dim3(1), dim3(kDataLen*2+10), N*N, stream, a, d_x, d_y);
axpy_2<float,double><<<1, kDataLen*2+10, N*N, stream>>>(a, d_x, d_y);
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_2<float,double>), dim3(1,1,1), dim3(kDataLen*2+10), N*N, stream, a, d_x, d_y);
axpy_2<float,double><<<dim3(1,1,1), kDataLen*2+10, N*N, stream>>>(a, d_x, d_y);
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_2<float,double>), dim3(1), dim3(kDataLen*2+10), N*N, stream, a, d_x, d_y);
axpy_2<float,double><<<1, dim3(kDataLen*2+10), N*N, stream>>>(a, d_x, d_y);
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_2<float,double>), dim3(1,1,1), dim3(kDataLen*2+10), N*N, stream, a, d_x, d_y);
axpy_2<float,double><<<dim3(1,1,1), dim3(kDataLen*2+10), N*N, stream>>>(a, d_x, d_y);
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_empty<float>), dim3(1), dim3(kDataLen), 0, 0);
axpy_empty<float><<<1, kDataLen>>>();
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_empty<float>), dim3(1), dim3(kDataLen), 0, 0);
axpy_empty<float><<<dim3(1), kDataLen>>>();
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_empty<float>), dim3(1), dim3(kDataLen), 0, 0);
axpy_empty<float><<<1, dim3(kDataLen)>>>();
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_empty<float>), dim3(1), dim3(kDataLen), 0, 0);
axpy_empty<float><<<dim3(1), dim3(kDataLen)>>>();
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_empty<float>), dim3(1), dim3(kDataLen), N, 0);
axpy_empty<float><<<1, kDataLen, N>>>();
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_empty<float>), dim3(1), dim3(kDataLen), N, 0);
axpy_empty<float><<<dim3(1), kDataLen, N>>>();
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_empty<float>), dim3(1), dim3(kDataLen), N, 0);
axpy_empty<float><<<1, dim3(kDataLen), N>>>();
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_empty<float>), dim3(1), dim3(kDataLen), N, 0);
axpy_empty<float><<<dim3(1), dim3(kDataLen), N>>>();
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_empty<float>), dim3(1), dim3(kDataLen), N, stream);
axpy_empty<float><<<1, kDataLen, N, stream>>>();
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_empty<float>), dim3(1), dim3(kDataLen), N, stream);
axpy_empty<float><<<dim3(1), kDataLen, N, stream>>>();
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_empty<float>), dim3(1), dim3(kDataLen), N, stream);
axpy_empty<float><<<1, dim3(kDataLen), N, stream>>>();
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_empty<float>), dim3(1), dim3(kDataLen), N, stream);
axpy_empty<float><<<dim3(1), dim3(kDataLen), N, stream>>>();
// CHECK: hipLaunchKernelGGL(empty, dim3(1), dim3(kDataLen), 0, 0);
empty<<<1, kDataLen>>> ( );
// CHECK: hipLaunchKernelGGL(empty, dim3(1), dim3(kDataLen), 0, 0);
empty<<<dim3(1), kDataLen>>> ( );
// CHECK: hipLaunchKernelGGL(empty, dim3(1), dim3(kDataLen), 0, 0);
empty<<<1, dim3(kDataLen)>>> ( );
// CHECK: hipLaunchKernelGGL(empty, dim3(1), dim3(kDataLen), 0, 0);
empty<<<dim3(1), dim3(kDataLen)>>> ( );
// CHECK: hipLaunchKernelGGL(empty, dim3(1), dim3(kDataLen), N, 0);
empty<<<1, kDataLen, N>>> ( );
// CHECK: hipLaunchKernelGGL(empty, dim3(1), dim3(kDataLen), N, 0);
empty<<<dim3(1), kDataLen, N>>> ( );
// CHECK: hipLaunchKernelGGL(empty, dim3(1), dim3(kDataLen), N, 0);
empty<<<1, dim3(kDataLen), N>>> ( );
// CHECK: hipLaunchKernelGGL(empty, dim3(1), dim3(kDataLen), N, 0);
empty<<<dim3(1), dim3(kDataLen), N>>> ( );
// CHECK: hipLaunchKernelGGL(empty, dim3(1), dim3(kDataLen), N, stream);
empty<<<1, kDataLen, N, stream>>> ( );
// CHECK: hipLaunchKernelGGL(empty, dim3(1), dim3(kDataLen), N, stream);
empty<<<dim3(1), kDataLen, N, stream>>> ( );
// CHECK: hipLaunchKernelGGL(empty, dim3(1), dim3(kDataLen), N, stream);
empty<<<1, dim3(kDataLen), N, stream>>> ( );
// CHECK: hipLaunchKernelGGL(empty, dim3(1), dim3(kDataLen), N, stream);
empty<<<dim3(1), dim3(kDataLen), N, stream>>> ( );
// CHECK: hipLaunchKernelGGL(nonempty, dim3(1), dim3(kDataLen), 0, 0, x, y, z);
nonempty<<<1, kDataLen>>> (x, y, z);
// CHECK: hipLaunchKernelGGL(nonempty, dim3(1), dim3(kDataLen), 0, 0, x, y, z);
nonempty<<<dim3(1), kDataLen>>> (x, y, z);
// CHECK: hipLaunchKernelGGL(nonempty, dim3(1), dim3(kDataLen), 0, 0, x, y, z);
nonempty<<<1, dim3(kDataLen)>>> (x, y, z);
// CHECK: hipLaunchKernelGGL(nonempty, dim3(1), dim3(kDataLen), 0, 0, x, y, z);
nonempty<<<dim3(1), dim3(kDataLen)>>> (x, y, z);
// CHECK: hipLaunchKernelGGL(nonempty, dim3(1), dim3(kDataLen), N, 0, x, y, z);
nonempty<<<1, kDataLen, N>>> (x, y, z);
// CHECK: hipLaunchKernelGGL(nonempty, dim3(1), dim3(kDataLen), N, 0, x, y, z);
nonempty<<<dim3(1), kDataLen, N>>> (x, y, z);
// CHECK: hipLaunchKernelGGL(nonempty, dim3(1), dim3(kDataLen), N, 0, x, y, z);
nonempty<<<1, dim3(kDataLen), N>>> (x, y, z);
// CHECK: hipLaunchKernelGGL(nonempty, dim3(1), dim3(kDataLen), N, 0, x, y, z);
nonempty<<<dim3(1), dim3(kDataLen), N>>> (x, y, z);
// CHECK: hipLaunchKernelGGL(nonempty, dim3(1), dim3(kDataLen), N, stream, x, y, z);
nonempty<<<1, kDataLen, N, stream>>> (x, y, z);
// CHECK: hipLaunchKernelGGL(nonempty, dim3(1), dim3(kDataLen), N, stream, x, y, z);
nonempty<<<dim3(1), kDataLen, N, stream>>> (x, y, z);
// CHECK: hipLaunchKernelGGL(nonempty, dim3(1), dim3(kDataLen), N, stream, x, y, z);
nonempty<<<1, dim3(kDataLen), N, stream>>> (x, y, z);
// CHECK: hipLaunchKernelGGL(nonempty, dim3(1), dim3(kDataLen), N, stream, x, y, z);
nonempty<<<dim3(1), dim3(kDataLen), N, stream>>> (x, y, z);
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_2<float,double>), dim3(x,y,z), dim3(std::min(kDataLen*2+10,x)), std::min(x,y), stream, a, std::min(d_x,d_y), std::max(d_x,d_y));
axpy_2<float,double><<<dim3(x,y,z), std::min(kDataLen*2+10,x), std::min(x,y), stream>>>(a, std::min(d_x,d_y), std::max(d_x,d_y));
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_2<float,double>), dim3(x,y,z), dim3(std::min(kDataLen*2+10,x)), std::min(x,y), 0, a, std::min(d_x,d_y), std::max(d_x,d_y));
axpy_2<float,double><<<dim3(x,y,z), std::min(kDataLen*2+10,x), std::min(x,y)>>>(a, std::min(d_x,d_y), std::max(d_x,d_y));
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_2<float,double>), dim3(x,y,z), dim3(std::min(kDataLen*2+10,x)), 0, 0, a, std::min(d_x,d_y), std::max(d_x,d_y));
axpy_2<float,double><<<dim3(x,y,z), std::min(kDataLen*2+10,x)>>>(a, std::min(d_x,d_y), std::max(d_x,d_y));
// CHECK: hipLaunchKernelGGL(nonempty, dim3(x,y,z), dim3(x,y,std::min(y,z)), 0, 0, x, y, z);
nonempty<<<dim3(x,y,z), dim3(x,y,std::min(y,z))>>>(x, y, z);
// CHECK: hipLaunchKernelGGL(nonempty, dim3(x,y,z), dim3(x,y,std::min(std::max(x,y),z)), 0, 0, x, y, z);
nonempty<<<dim3(x,y,z), dim3(x,y,std::min(std::max(x,y),z))>>>(x, y, z);
// CHECK: hipLaunchKernelGGL(nonempty, dim3(x,y,z), dim3(x,y,std::min(std::max(x,int(N)),z)), 0, 0, x, y, z);
nonempty<<<dim3(x,y,z), dim3(x,y,std::min(std::max(x,int(N)),z))>>>(x, y, z);
// CHECK: hipLaunchKernelGGL(nonempty, dim3(x,y,z), dim3(x,y,std::min(std::max(x,int(N+N -x/y + y*1)),z)), 0, 0, x, y, z);
nonempty<<<dim3(x,y,z), dim3(x,y,std::min(std::max(x,int(N+N -x/y + y*1)),z))>>>(x, y, z);
// Copy output data to host.
// CHECK: hipDeviceSynchronize();
cudaDeviceSynchronize();
// CHECK: hipMemcpy(host_y, device_y, kDataLen * sizeof(float), hipMemcpyDeviceToHost);
cudaMemcpy(host_y, device_y, kDataLen * sizeof(float), cudaMemcpyDeviceToHost);
// Print the results.
for (int i = 0; i < kDataLen; ++i) {
std::cout << "y[" << i << "] = " << host_y[i] << "\n";
}
// CHECK: hipDeviceReset();
cudaDeviceReset();
return 0;
}
|
