Change the #define of GENERIC_GRID_LAUNCH to take valueat compilation, disable warning messages
Change-Id: Ic6c011529e26de359bcda1e7083727e7ee52887b
Этот коммит содержится в:
@@ -196,7 +196,7 @@ namespace glo_tests
|
||||
template<typename... Ts>
|
||||
void operator()(Ts&&...) const {}
|
||||
};
|
||||
#warning "GGL hipLaunchKernel defined"
|
||||
//#warning "GGL hipLaunchKernel defined"
|
||||
#define hipLaunchKernel( \
|
||||
kernel_name, \
|
||||
num_blocks, \
|
||||
|
||||
@@ -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* );
|
||||
|
||||
@@ -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, ...) \
|
||||
|
||||
@@ -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 <stdint.h>
|
||||
#include <stddef.h>
|
||||
#include <iostream>
|
||||
|
||||
#ifndef GENERIC_GRID_LAUNCH
|
||||
#define GENERIC_GRID_LAUNCH 0
|
||||
#endif
|
||||
|
||||
#include <hip/hcc_detail/host_defines.h>
|
||||
#include <hip/hip_runtime_api.h>
|
||||
#include <hip/hip_texture.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
|
||||
|
||||
|
||||
@@ -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 <>
|
||||
|
||||
Ссылка в новой задаче
Block a user