AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[HIPIFY][#207][fix] Translate all preprocessor's conditional blocks

Browse Source 
+ Start to translate preprocessor's false conditional blocks too:
  based on clang's https://reviews.llvm.org/D66597;
  available only starting from LLVM 10.0 or trunk.
+ Option -skip-excluded-preprocessor-conditional-blocks for skipping excluded conditional blocks:
  the default behavior for hipify-clang built with LLVM < 10.0;
  false by default for hipify-clang built with LLVM 10 or trunk.
+ Add 4 preprocessor unit tests, 2 of which are LLVM 10.0 only
+ Update couple of existing tests by setting -skip-excluded-preprocessor-conditional-blocks option:
  update lit testing accordingly
This commit is contained in:
Evgeny Mankov
2019-08-28 21:17:35 +03:00
parent 1ab81d9ce6
commit 24be21495d
15 changed files with 223 additions and 12 deletions
Download Patch File Download Diff File Expand all files Collapse all files
tests/hipify-clang/lit.cfg
+4
View File
@@ -42,6 +42,10 @@ if config.cuda_version_major < 10:
config.excludes.append('cuSPARSE_10.cu')
config.excludes.append('cuSPARSE_11.cu')
if config.llvm_version_major < 10:
config.excludes.append('pp_if_else_conditionals_LLVM_10.cu')
config.excludes.append('pp_if_else_conditionals_01_LLVM_10.cu')
# name: The name of this test suite.
config.name = 'hipify'
tests/hipify-clang/lit.site.cfg.in
+1
View File
@@ -3,6 +3,7 @@ import os
config.pointer_size = @CMAKE_SIZEOF_VOID_P@
config.llvm_version = "@LLVM_PACKAGE_VERSION@"
config.llvm_version_major = int("@LLVM_VERSION_MAJOR@")
config.llvm_tools_dir = "@LLVM_TOOLS_BINARY_DIR@"
config.obj_root = "@CMAKE_CURRENT_BINARY_DIR@"
config.cuda_root = "@CUDA_TOOLKIT_ROOT_DIR@"
tests/hipify-clang/unit_tests/libraries/cuComplex/cuComplex_Julia.cu
+1 -1
View File
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// RUN: %run_test hipify "%s" "%t" %hipify_args "--skip-excluded-preprocessor-conditional-blocks" %clang_args
// CHECK: #include <hip/hip_runtime.h>
// CHECK: #include "hip/hip_complex.h"
tests/hipify-clang/unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu
+1 -1
View File
@@ -1,4 +1,4 @@
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// RUN: %run_test hipify "%s" "%t" %hipify_args "--skip-excluded-preprocessor-conditional-blocks" %clang_args
#include <stdio.h>
#include <stdlib.h>
#include <assert.h>
tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals.cu
+30
View File
@@ -0,0 +1,30 @@
// RUN: %run_test hipify "%s" "%t" %hipify_args "--skip-excluded-preprocessor-conditional-blocks" %clang_args
// CHECK: #include <hip/hip_runtime.h>
#include <cuda.h>
__global__ void axpy_kernel(float a, float* x, float* y) {
y[threadIdx.x] = a * x[threadIdx.x];
}
void axpy(float a, float* x, float* y) {
#ifdef SOME_MACRO
// CHECK: axpy_kernel <<<1, 1>>> (a, y, x);
axpy_kernel <<<1, 1>>> (a, y, x);
#endif
#ifndef SOME_MACRO
// CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(2), 0, 0, a, y, x);
axpy_kernel <<<1, 2>>> (a, y, x);
#endif
#ifdef SOME_MACRO
// CHECK: axpy_kernel <<<1, 3>>> (a, y, x);
axpy_kernel <<<1, 3>>> (a, y, x);
#else
// CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(4), 0, 0, a, x, y);
axpy_kernel <<<1, 4>>> (a, x, y);
#endif
}
tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals_01.cu
+52
View File
@@ -0,0 +1,52 @@
// RUN: %run_test hipify "%s" "%t" %hipify_args "--skip-excluded-preprocessor-conditional-blocks" %clang_args
// CHECK: #include <hip/hip_runtime.h>
__global__ void axpy_kernel(float a, float* x, float* y) {
y[threadIdx.x] = a * x[threadIdx.x];
}
void axpy(float a, float* x, float* y) {
float* y_new = nullptr;
#ifdef SOME_MACRO
y_new = x;
// CHECK: axpy_kernel <<<1, 1>>> (a, y_new, x);
axpy_kernel <<<1, 1>>> (a, y_new, x);
#endif
#ifndef SOME_MACRO
y_new = y;
// CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(2), 0, 0, a, y_new, x);
axpy_kernel <<<1, 2>>> (a, y_new, x);
#endif
#ifdef SOME_MACRO
// CHECK: axpy_kernel <<<1, 3>>> (a, y, x);
axpy_kernel <<<1, 3>>> (a, y, x);
#else
// CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(4), 0, 0, a, x, y);
axpy_kernel <<<1, 4>>> (a, x, y);
#endif
#ifdef SOME_MACRO
// CHECK: axpy_kernel <<<1, 5>>> (a, y, x);
axpy_kernel <<<1, 5>>> (a, y, x);
#elif defined SOME_MACRO_1
// CHECK: axpy_kernel <<<1, 6>>> (a, x, y);
axpy_kernel <<<1, 6>>> (a, x, y);
#else
// CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(7), 0, 0, a, x, y);
axpy_kernel <<<1, 7>>> (a, x, y);
#endif
#ifndef SOME_MACRO
// CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(8), 0, 0, a, y, x);
axpy_kernel <<<1, 8>>> (a, y, x);
#elif !defined(SOME_MACRO_1)
// CHECK: axpy_kernel <<<1, 9>>> (a, x, y);
axpy_kernel <<<1, 9>>> (a, x, y);
#else
// CHECK: axpy_kernel <<<1, 10>>> (a, x, y);
axpy_kernel <<<1, 10>>> (a, x, y);
#endif
}
tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals_01_LLVM_10.cu
+52
View File
@@ -0,0 +1,52 @@
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// CHECK: #include <hip/hip_runtime.h>
__global__ void axpy_kernel(float a, float* x, float* y) {
y[threadIdx.x] = a * x[threadIdx.x];
}
void axpy(float a, float* x, float* y) {
float* y_new = nullptr;
#ifdef SOME_MACRO
y_new = x;
// CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(1), 0, 0, a, y_new, x);
axpy_kernel <<<1, 1>>> (a, y_new, x);
#endif
#ifndef SOME_MACRO
y_new = y;
// CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(2), 0, 0, a, y_new, x);
axpy_kernel <<<1, 2>>> (a, y_new, x);
#endif
#ifdef SOME_MACRO
// CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(3), 0, 0, a, y, x);
axpy_kernel <<<1, 3>>> (a, y, x);
#else
// CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(4), 0, 0, a, x, y);
axpy_kernel <<<1, 4>>> (a, x, y);
#endif
#ifdef SOME_MACRO
// CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(5), 0, 0, a, y, x);
axpy_kernel <<<1, 5>>> (a, y, x);
#elif defined SOME_MACRO_1
// CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(6), 0, 0, a, x, y);
axpy_kernel <<<1, 6>>> (a, x, y);
#else
// CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(7), 0, 0, a, x, y);
axpy_kernel <<<1, 7>>> (a, x, y);
#endif
#ifndef SOME_MACRO
// CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(8), 0, 0, a, y, x);
axpy_kernel <<<1, 8>>> (a, y, x);
#elif !defined(SOME_MACRO_1)
// CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(9), 0, 0, a, x, y);
axpy_kernel <<<1, 9>>> (a, x, y);
#else
// CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(10), 0, 0, a, x, y);
axpy_kernel <<<1, 10>>> (a, x, y);
#endif
}
tests/hipify-clang/unit_tests/pp/pp_if_else_conditionals_LLVM_10.cu
+30
View File
@@ -0,0 +1,30 @@
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// CHECK: #include <hip/hip_runtime.h>
#include <cuda.h>
__global__ void axpy_kernel(float a, float* x, float* y) {
y[threadIdx.x] = a * x[threadIdx.x];
}
void axpy(float a, float* x, float* y) {
#ifdef SOME_MACRO
// CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(1), 0, 0, a, y, x);
axpy_kernel <<<1, 1>>> (a, y, x);
#endif
#ifndef SOME_MACRO
// CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(2), 0, 0, a, y, x);
axpy_kernel <<<1, 2>>> (a, y, x);
#endif
#ifdef SOME_MACRO
// CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(3), 0, 0, a, y, x);
axpy_kernel <<<1, 3>>> (a, y, x);
#else
// CHECK: hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(4), 0, 0, a, x, y);
axpy_kernel <<<1, 4>>> (a, x, y);
#endif
}