Chris Kitching
6491c2c3eb
All but the most contrived use of macros is now properly handled - have a look at the new testcases this commit adds. You can have macros in kernel calls, macros spanning chunks of your arguments, the call, call parameters, or callee can all be macros or partially macros.
81 lines
2.7 KiB
Plaintext
81 lines
2.7 KiB
Plaintext
// RUN: %run_test hipify "%s" "%t" %cuda_args
|
|
|
|
#include <iostream>
|
|
|
|
|
|
#define TOKEN_PASTE(X, Y) X ## Y
|
|
#define ARG_LIST_AS_MACRO a, device_x, device_y
|
|
#define KERNEL_CALL_AS_MACRO axpy<float><<<1, kDataLen>>>
|
|
#define KERNEL_NAME_MACRO axpy<float>
|
|
|
|
// CHECK: #define COMPLETE_LAUNCH hipLaunchKernelGGL(axpy, dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y)
|
|
#define COMPLETE_LAUNCH axpy<<<1, kDataLen>>>(a, device_x, device_y)
|
|
|
|
|
|
template<typename T>
|
|
__global__ void axpy(T a, T *x, T *y) {
|
|
// CHECK: y[hipThreadIdx_x] = a * x[hipThreadIdx_x];
|
|
y[threadIdx.x] = a * x[threadIdx.x];
|
|
}
|
|
|
|
|
|
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);
|
|
|
|
// Launch the kernel in numerous different strange ways to exercise the prerocessor.
|
|
// CHECK: hipLaunchKernelGGL(axpy, dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y);
|
|
axpy<<<1, kDataLen>>>(a, device_x, device_y);
|
|
|
|
// CHECK: hipLaunchKernelGGL(axpy<float>, dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y);
|
|
axpy<float><<<1, kDataLen>>>(a, device_x, device_y);
|
|
|
|
// CHECK: hipLaunchKernelGGL(axpy<float>, dim3(1), dim3(kDataLen), 0, 0, a, TOKEN_PASTE(device, _x), device_y);
|
|
axpy<float><<<1, kDataLen>>>(a, TOKEN_PASTE(device, _x), device_y);
|
|
|
|
// CHECK: hipLaunchKernelGGL(axpy<float>, dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO);
|
|
axpy<float><<<1, kDataLen>>>(ARG_LIST_AS_MACRO);
|
|
|
|
// CHECK: hipLaunchKernelGGL(KERNEL_NAME_MACRO, dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO);
|
|
KERNEL_NAME_MACRO<<<1, kDataLen>>>(ARG_LIST_AS_MACRO);
|
|
|
|
// CHECK: hipLaunchKernelGGL(axpy<float>, dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO);
|
|
KERNEL_CALL_AS_MACRO(ARG_LIST_AS_MACRO);
|
|
|
|
// CHECK: COMPLETE_LAUNCH;
|
|
COMPLETE_LAUNCH;
|
|
|
|
|
|
// 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;
|
|
}
|