From faf0fbebe9d21671435f4c24be076fd4445ee945 Mon Sep 17 00:00:00 2001 From: pensun Date: Fri, 17 Mar 2017 19:57:20 +0000 Subject: [PATCH] Change the #define of GENERIC_GRID_LAUNCH to take valueat compilation, disable warning messages Change-Id: Ic6c011529e26de359bcda1e7083727e7ee52887b --- hipamd/include/hip/hcc_detail/grid_launch_v2.hpp | 2 +- hipamd/include/hip/hcc_detail/hip_ldg.h | 1 - hipamd/include/hip/hcc_detail/hip_runtime.h | 8 ++++---- hipamd/include/hip/hcc_detail/hip_runtime_api.h | 6 ++++-- hipamd/include/hip/hcc_detail/host_defines.h | 13 +++++++++---- hipamd/src/trace_helper.h | 2 -- 6 files changed, 18 insertions(+), 14 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/grid_launch_v2.hpp b/hipamd/include/hip/hcc_detail/grid_launch_v2.hpp index ab11433a5b..9f7492b4d5 100644 --- a/hipamd/include/hip/hcc_detail/grid_launch_v2.hpp +++ b/hipamd/include/hip/hcc_detail/grid_launch_v2.hpp @@ -196,7 +196,7 @@ namespace glo_tests template void operator()(Ts&&...) const {} }; -#warning "GGL hipLaunchKernel defined" +//#warning "GGL hipLaunchKernel defined" #define hipLaunchKernel( \ kernel_name, \ num_blocks, \ diff --git a/hipamd/include/hip/hcc_detail/hip_ldg.h b/hipamd/include/hip/hcc_detail/hip_ldg.h index 6bf7a618d0..5c33ee773f 100644 --- a/hipamd/include/hip/hcc_detail/hip_ldg.h +++ b/hipamd/include/hip/hcc_detail/hip_ldg.h @@ -27,7 +27,6 @@ THE SOFTWARE. #if __hcc_workweek__ >= 16164 #include "hip_vector_types.h" #include "host_defines.h" -#warning "LDG header included" __device__ char __ldg(const char* ); __device__ char2 __ldg(const char2* ); diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h index 9d0a43b1df..673807fcd2 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -63,20 +63,20 @@ THE SOFTWARE. //TODO-HCC-GL - change this to typedef. //typedef grid_launch_parm hipLaunchParm ; struct EmptyLaunchParm{}; -#ifndef GENERIC_GRID_LAUNCH +#if GENERIC_GRID_LAUNCH == 0 #define hipLaunchParm grid_launch_parm #else #define hipLaunchParm EmptyLaunchParm #endif //GENERIC_GRID_LAUNCH -#if defined (GRID_LAUNCH_VERSION) and (GRID_LAUNCH_VERSION >= 20) || defined GENERIC_GRID_LAUNCH +#if defined (GRID_LAUNCH_VERSION) and (GRID_LAUNCH_VERSION >= 20) || GENERIC_GRID_LAUNCH == 1 #else // Use field names for grid_launch 2.0 structure, if HCC supports GL 2.0. #error (HCC must support GRID_LAUNCH_20) #endif //GRID_LAUNCH_VERSION #endif //HCC -#if defined GENERIC_GRID_LAUNCH && defined __HCC__ +#if GENERIC_GRID_LAUNCH==1 && defined __HCC__ #include "grid_launch_v2.hpp" #endif//GENERIC_GRID_LAUNCH @@ -426,7 +426,7 @@ extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 blo extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, grid_launch_parm *lp, const char *kernelNameStr); extern void ihipPostLaunchKernel(const char *kernelName, hipStream_t stream, grid_launch_parm &lp); -#ifndef GENERIC_GRID_LAUNCH +#if GENERIC_GRID_LAUNCH == 0 #warning "Original hipLaunchKernel defined" // Due to multiple overloaded versions of ihipPreLaunchKernel, the numBlocks3D and blockDim3D can be either size_t or dim3 types #define hipLaunchKernel(_kernelName, _numBlocks3D, _blockDim3D, _groupMemBytes, _stream, ...) \ diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index c769548214..01fea548b8 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -27,12 +27,14 @@ THE SOFTWARE. * @file hcc_detail/hip_runtime_api.h * @brief Contains C function APIs for HIP runtime. This file does not use any HCC builtin or special language extensions (-hc mode) ; those functions in hip_runtime.h. */ -// guard for grid_launch_v2 -#define GENERIC_GRID_LAUNCH #include #include #include +#ifndef GENERIC_GRID_LAUNCH +#define GENERIC_GRID_LAUNCH 0 +#endif + #include #include #include diff --git a/hipamd/include/hip/hcc_detail/host_defines.h b/hipamd/include/hip/hcc_detail/host_defines.h index d7128d6fab..6804ba464b 100644 --- a/hipamd/include/hip/hcc_detail/host_defines.h +++ b/hipamd/include/hip/hcc_detail/host_defines.h @@ -30,6 +30,11 @@ THE SOFTWARE. #define USE_PROMOTE_FREE_HCC 0 +// Add guard to Generic Grid Launch method +#ifndef GENERIC_GRID_LAUNCH +#define GENERIC_GRID_LAUNCH 0 +#endif + #ifdef __HCC__ /** * Function and kernel markers @@ -37,12 +42,12 @@ THE SOFTWARE. #define __host__ __attribute__((cpu)) #define __device__ __attribute__((hc)) -#warning "HOST DEFINE header included" -#ifndef GENERIC_GRID_LAUNCH -#warning "original global define reached" +//#warning "HOST DEFINE header included" +#if GENERIC_GRID_LAUNCH == 0 +//#warning "original global define reached" #define __global__ __attribute__((hc_grid_launch)) __attribute__((used)) #else -#warning "GGL global define reached" +//#warning "GGL global define reached" #define __global__ [[hc]] __attribute__((weak)) #endif //GENERIC_GRID_LAUNCH diff --git a/hipamd/src/trace_helper.h b/hipamd/src/trace_helper.h index abff491916..d49cb67be0 100644 --- a/hipamd/src/trace_helper.h +++ b/hipamd/src/trace_helper.h @@ -70,7 +70,6 @@ inline std::string ToString(hipEvent_t v) ss << v; return ss.str(); }; -#ifndef GENERIC_GRID_LAUNCH // hipStream_t template <> inline std::string ToString(hipStream_t v) @@ -84,7 +83,6 @@ inline std::string ToString(hipStream_t v) return ss.str(); }; -#endif //GENERIC_GRID_LAUNCH // hipMemcpyKind specialization template <>