diff --git a/hipamd/CMakeLists.txt b/hipamd/CMakeLists.txt index 0c209b8fa2..18e8e7df4d 100644 --- a/hipamd/CMakeLists.txt +++ b/hipamd/CMakeLists.txt @@ -53,6 +53,7 @@ if(${HIPIFY_STANDLONE}) set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${EXTRA_CFLAGS}") set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${EXTRA_CFLAGS}") set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wall -std=c++11 -pthread -fno-rtti -fvisibility-inlines-hidden") + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DHIPIFY_CLANG_RES=\\\"${LLVM_LIBRARY_DIRS}/clang/${LLVM_VERSION_MAJOR}.${LLVM_VERSION_MINOR}.${LLVM_VERSION_PATCH}\\\"") add_subdirectory(test) else() install(TARGETS hipify diff --git a/hipamd/src/Cuda2Hip.cpp b/hipamd/src/Cuda2Hip.cpp index 434d74e584..a9cb162849 100644 --- a/hipamd/src/Cuda2Hip.cpp +++ b/hipamd/src/Cuda2Hip.cpp @@ -685,6 +685,9 @@ int main(int argc, const char **argv) { { Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster(Stage, ArgumentInsertPosition::BEGIN)); Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster("-std=c++11")); +#if defined(HIPIFY_CLANG_RES) + Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES)); +#endif // defined(HIPIFY_CLANG_HEADERS) Tool.appendArgumentsAdjuster(getClangSyntaxOnlyAdjuster()); Result = Tool.run(action.get()); diff --git a/hipamd/test/axpy.cu b/hipamd/test/axpy.cu index 29d5493763..9e83ccb7e6 100644 --- a/hipamd/test/axpy.cu +++ b/hipamd/test/axpy.cu @@ -1,11 +1,10 @@ -// RUN: hipify "%s" 2>&1 | FileCheck %s - -#include // for checkCudaErrors +// RUN: hipify "%s" -o=%t -- +// RUN: FileCheck %s -input-file=%t --match-full-lines #include __global__ void axpy(float a, float* x, float* y) { - // CHECK: hipThreadIdx_x + // CHECK: y[hipThreadIdx_x] = a * x[hipThreadIdx_x]; y[threadIdx.x] = a * x[threadIdx.x]; } @@ -19,25 +18,23 @@ int main(int argc, char* argv[]) { // Copy input data to device. float* device_x; float* device_y; - checkCudaErrors(cudaMalloc(&device_x, kDataLen * sizeof(float))); - checkCudaErrors(cudaMalloc(&device_y, kDataLen * sizeof(float))); - checkCudaErrors(cudaMemcpy(device_x, host_x, kDataLen * sizeof(float), - cudaMemcpyHostToDevice)); + cudaMalloc(&device_x, kDataLen * sizeof(float)); + cudaMalloc(&device_y, kDataLen * sizeof(float)); + cudaMemcpy(device_x, host_x, kDataLen * sizeof(float), cudaMemcpyHostToDevice); // Launch the kernel. - // CHECK: hipLaunchKernel(HIP_KERNEL_NAME + // CHECK: hipLaunchKernel(HIP_KERNEL_NAME(axpy), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y); axpy<<<1, kDataLen>>>(a, device_x, device_y); // Copy output data to host. - checkCudaErrors(cudaDeviceSynchronize()); - checkCudaErrors(cudaMemcpy(host_y, device_y, kDataLen * sizeof(float), - cudaMemcpyDeviceToHost)); + cudaDeviceSynchronize(); + 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"; } - checkCudaErrors(cudaDeviceReset()); + cudaDeviceReset(); return 0; }