29 #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_RUNTIME_H 30 #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_RUNTIME_H 32 #include <hip/hcc_detail/hip_common.h> 46 #if __HCC_OR_HIP_CLANG__ 49 #if !defined(__align__) 50 #define __align__(x) __attribute__((aligned(x))) 54 #define CUDA_SUCCESS hipSuccess 56 #include <hip/hip_runtime_api.h> 57 #endif // __HCC_OR_HIP_CLANG__ 61 #ifdef HIP_ENABLE_PRINTF 62 #define HCC_ENABLE_ACCELERATOR_PRINTF 1 68 #include "grid_launch.h" 69 #include "hc_printf.hpp" 73 #if GENERIC_GRID_LAUNCH == 0 74 #define hipLaunchParm grid_launch_parm 77 struct Empty_launch_parm {};
79 #define hipLaunchParm hip_impl::Empty_launch_parm 80 #endif // GENERIC_GRID_LAUNCH 82 #if defined(GRID_LAUNCH_VERSION) and (GRID_LAUNCH_VERSION >= 20) || GENERIC_GRID_LAUNCH == 1 83 #else // Use field names for grid_launch 2.0 structure, if HCC supports GL 2.0. 84 #error(HCC must support GRID_LAUNCH_20) 85 #endif // GRID_LAUNCH_VERSION 89 #if GENERIC_GRID_LAUNCH == 1 && defined __HCC__ 90 #include "grid_launch_GGL.hpp" 91 #endif // GENERIC_GRID_LAUNCH 95 #if __HCC_OR_HIP_CLANG__ 96 extern int HIP_TRACE_API;
99 #include <hip/hcc_detail/hip_ldg.h> 101 #include <hip/hcc_detail/hip_atomic.h> 103 #include <hip/hcc_detail/device_functions.h> 104 #include <hip/hcc_detail/surface_functions.h> 105 #include <hip/hcc_detail/texture_functions.h> 107 #include <hip/hcc_detail/math_functions.h> 110 #if defined(__KALMAR_ACCELERATOR__) && !defined(__HCC_ACCELERATOR__) 111 #define __HCC_ACCELERATOR__ __KALMAR_ACCELERATOR__ 115 #if __HIP_DEVICE_COMPILE__ == 1 117 #define assert(COND) \ 127 #if (defined(__HCC_ACCELERATOR__) && (__HCC_ACCELERATOR__ != 0)) || __HIP_DEVICE_COMPILE__ 131 #define __HIP_ARCH_HAS_GLOBAL_INT32_ATOMICS__ (1) 132 #define __HIP_ARCH_HAS_GLOBAL_FLOAT_ATOMIC_EXCH__ (1) 133 #define __HIP_ARCH_HAS_SHARED_INT32_ATOMICS__ (1) 134 #define __HIP_ARCH_HAS_SHARED_FLOAT_ATOMIC_EXCH__ (1) 135 #define __HIP_ARCH_HAS_FLOAT_ATOMIC_ADD__ (1) 138 #define __HIP_ARCH_HAS_GLOBAL_INT64_ATOMICS__ (1) 139 #define __HIP_ARCH_HAS_SHARED_INT64_ATOMICS__ (0) 142 #define __HIP_ARCH_HAS_DOUBLES__ (1) 145 #define __HIP_ARCH_HAS_WARP_VOTE__ (1) 146 #define __HIP_ARCH_HAS_WARP_BALLOT__ (1) 147 #define __HIP_ARCH_HAS_WARP_SHUFFLE__ (1) 148 #define __HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__ (0) 151 #define __HIP_ARCH_HAS_THREAD_FENCE_SYSTEM__ (1) 152 #define __HIP_ARCH_HAS_SYNC_THREAD_EXT__ (0) 155 #define __HIP_ARCH_HAS_SURFACE_FUNCS__ (0) 156 #define __HIP_ARCH_HAS_3DGRID__ (1) 157 #define __HIP_ARCH_HAS_DYNAMIC_PARALLEL__ (0) 162 #define launch_bounds_impl0(requiredMaxThreadsPerBlock) \ 163 __attribute__((amdgpu_flat_work_group_size(1, requiredMaxThreadsPerBlock))) 164 #define launch_bounds_impl1(requiredMaxThreadsPerBlock, minBlocksPerMultiprocessor) \ 165 __attribute__((amdgpu_flat_work_group_size(1, requiredMaxThreadsPerBlock), \ 166 amdgpu_waves_per_eu(minBlocksPerMultiprocessor))) 167 #define select_impl_(_1, _2, impl_, ...) impl_ 168 #define __launch_bounds__(...) \ 169 select_impl_(__VA_ARGS__, launch_bounds_impl1, launch_bounds_impl0)(__VA_ARGS__) 172 #if defined(__cplusplus) 174 #elif defined(__STDC_VERSION__) 178 __host__ inline void* __get_dynamicgroupbaseptr() {
return nullptr; }
180 #if __HIP_ARCH_GFX701__ == 0 182 __device__
unsigned __hip_ds_bpermute(
int index,
unsigned src);
183 __device__
float __hip_ds_bpermutef(
int index,
float src);
184 __device__
unsigned __hip_ds_permute(
int index,
unsigned src);
185 __device__
float __hip_ds_permutef(
int index,
float src);
187 template <
int pattern>
188 __device__
unsigned __hip_ds_swizzle_N(
unsigned int src);
189 template <
int pattern>
190 __device__
float __hip_ds_swizzlef_N(
float src);
192 template <
int dpp_ctrl,
int row_mask,
int bank_mask,
bool bound_ctrl>
193 __device__
int __hip_move_dpp_N(
int src);
195 #endif //__HIP_ARCH_GFX803__ == 1 197 #endif // __HCC_OR_HIP_CLANG__ 202 typename std::common_type<decltype(hc_get_group_id), decltype(hc_get_group_size),
203 decltype(hc_get_num_groups), decltype(hc_get_workitem_id)>::type f>
205 using R = decltype(f(0));
208 __device__
operator R()
const {
return f(0); }
211 __device__
operator R()
const {
return f(1); }
214 __device__
operator R()
const {
return f(2); }
218 static constexpr X x{};
219 static constexpr Y y{};
220 static constexpr Z z{};
228 #define hipThreadIdx_x (hc_get_workitem_id(0)) 229 #define hipThreadIdx_y (hc_get_workitem_id(1)) 230 #define hipThreadIdx_z (hc_get_workitem_id(2)) 232 #define hipBlockIdx_x (hc_get_group_id(0)) 233 #define hipBlockIdx_y (hc_get_group_id(1)) 234 #define hipBlockIdx_z (hc_get_group_id(2)) 236 #define hipBlockDim_x (hc_get_group_size(0)) 237 #define hipBlockDim_y (hc_get_group_size(1)) 238 #define hipBlockDim_z (hc_get_group_size(2)) 240 #define hipGridDim_x (hc_get_num_groups(0)) 241 #define hipGridDim_y (hc_get_num_groups(1)) 242 #define hipGridDim_z (hc_get_num_groups(2)) 244 #endif // defined __HCC__ 245 #if __HCC_OR_HIP_CLANG__ 246 extern "C" __device__
void* __hip_malloc(
size_t);
247 extern "C" __device__
void* __hip_free(
void* ptr);
249 static inline __device__
void* malloc(
size_t size) {
return __hip_malloc(size); }
250 static inline __device__
void* free(
void* ptr) {
return __hip_free(ptr); }
252 #if defined(__HCC_ACCELERATOR__) && defined(HC_FEATURE_PRINTF) 253 template <
typename... All>
254 static inline __device__
void printf(
const char* format, All... all) {
255 hc::printf(format, all...);
257 #elif defined(__HCC_ACCELERATOR__) || __HIP__ 258 template <
typename... All>
259 static inline __device__
void printf(
const char* format, All... all) {}
262 #endif //__HCC_OR_HIP_CLANG__ 266 #define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE) 268 #define HIP_KERNEL_NAME(...) (__VA_ARGS__) 269 #define HIP_SYMBOL(X) #X 271 #if defined __HCC_CPP__ 282 #if GENERIC_GRID_LAUNCH == 0 286 #define hipLaunchKernel(_kernelName, _numBlocks3D, _blockDim3D, _groupMemBytes, _stream, ...) \ 288 grid_launch_parm lp; \ 289 lp.dynamic_group_mem_bytes = _groupMemBytes; \ 290 hipStream_t trueStream = \ 291 (ihipPreLaunchKernel(_stream, _numBlocks3D, _blockDim3D, &lp, #_kernelName)); \ 292 _kernelName(lp, ##__VA_ARGS__); \ 293 ihipPostLaunchKernel(#_kernelName, trueStream, lp); \ 295 #endif // GENERIC_GRID_LAUNCH 297 #elif defined(__HCC_C__) 324 #elif defined(__clang__) && defined(__HIP__) 326 #define HIP_KERNEL_NAME(...) __VA_ARGS__ 327 #define HIP_SYMBOL(X) #X 329 typedef int hipLaunchParm;
331 #define hipLaunchKernel(kernelName, numblocks, numthreads, memperblock, streamId, ...) \ 333 kernelName<<<(numblocks), (numthreads), (memperblock), (streamId)>>>(hipLaunchParm{}, ##__VA_ARGS__); \ 336 #define hipLaunchKernelGGL(kernelName, numblocks, numthreads, memperblock, streamId, ...) \ 338 kernelName<<<(numblocks), (numthreads), (memperblock), (streamId)>>>(__VA_ARGS__); \ 341 #include <hip/hip_runtime_api.h> 343 #pragma push_macro("__DEVICE__") 344 #define __DEVICE__ static __device__ __forceinline__ 346 extern "C" __device__
size_t __ockl_get_local_id(uint);
347 __DEVICE__ uint __hip_get_thread_idx_x() {
return __ockl_get_local_id(0); }
348 __DEVICE__ uint __hip_get_thread_idx_y() {
return __ockl_get_local_id(1); }
349 __DEVICE__ uint __hip_get_thread_idx_z() {
return __ockl_get_local_id(2); }
351 extern "C" __device__
size_t __ockl_get_group_id(uint);
352 __DEVICE__ uint __hip_get_block_idx_x() {
return __ockl_get_group_id(0); }
353 __DEVICE__ uint __hip_get_block_idx_y() {
return __ockl_get_group_id(1); }
354 __DEVICE__ uint __hip_get_block_idx_z() {
return __ockl_get_group_id(2); }
356 extern "C" __device__
size_t __ockl_get_local_size(uint);
357 __DEVICE__ uint __hip_get_block_dim_x() {
return __ockl_get_local_size(0); }
358 __DEVICE__ uint __hip_get_block_dim_y() {
return __ockl_get_local_size(1); }
359 __DEVICE__ uint __hip_get_block_dim_z() {
return __ockl_get_local_size(2); }
361 extern "C" __device__
size_t __ockl_get_num_groups(uint);
362 __DEVICE__ uint __hip_get_grid_dim_x() {
return __ockl_get_num_groups(0); }
363 __DEVICE__ uint __hip_get_grid_dim_y() {
return __ockl_get_num_groups(1); }
364 __DEVICE__ uint __hip_get_grid_dim_z() {
return __ockl_get_num_groups(2); }
366 #define __HIP_DEVICE_BUILTIN(DIMENSION, FUNCTION) \ 367 __declspec(property(get = __get_##DIMENSION)) uint DIMENSION; \ 368 __DEVICE__ uint __get_##DIMENSION(void) { \ 372 struct __hip_builtin_threadIdx_t {
373 __HIP_DEVICE_BUILTIN(x,__hip_get_thread_idx_x());
374 __HIP_DEVICE_BUILTIN(y,__hip_get_thread_idx_y());
375 __HIP_DEVICE_BUILTIN(z,__hip_get_thread_idx_z());
378 struct __hip_builtin_blockIdx_t {
379 __HIP_DEVICE_BUILTIN(x,__hip_get_block_idx_x());
380 __HIP_DEVICE_BUILTIN(y,__hip_get_block_idx_y());
381 __HIP_DEVICE_BUILTIN(z,__hip_get_block_idx_z());
384 struct __hip_builtin_blockDim_t {
385 __HIP_DEVICE_BUILTIN(x,__hip_get_block_dim_x());
386 __HIP_DEVICE_BUILTIN(y,__hip_get_block_dim_y());
387 __HIP_DEVICE_BUILTIN(z,__hip_get_block_dim_z());
390 struct __hip_builtin_gridDim_t {
391 __HIP_DEVICE_BUILTIN(x,__hip_get_grid_dim_x());
392 __HIP_DEVICE_BUILTIN(y,__hip_get_grid_dim_y());
393 __HIP_DEVICE_BUILTIN(z,__hip_get_grid_dim_z());
396 #undef __HIP_DEVICE_BUILTIN 397 #pragma pop_macro("__DEVICE__") 399 extern const __device__
__attribute__((weak)) __hip_builtin_threadIdx_t threadIdx;
400 extern const __device__
__attribute__((weak)) __hip_builtin_blockIdx_t blockIdx;
401 extern const __device__
__attribute__((weak)) __hip_builtin_blockDim_t blockDim;
402 extern const __device__
__attribute__((weak)) __hip_builtin_gridDim_t gridDim;
405 #define hipThreadIdx_x threadIdx.x 406 #define hipThreadIdx_y threadIdx.y 407 #define hipThreadIdx_z threadIdx.z 409 #define hipBlockIdx_x blockIdx.x 410 #define hipBlockIdx_y blockIdx.y 411 #define hipBlockIdx_z blockIdx.z 413 #define hipBlockDim_x blockDim.x 414 #define hipBlockDim_y blockDim.y 415 #define hipBlockDim_z blockDim.z 417 #define hipGridDim_x gridDim.x 418 #define hipGridDim_y gridDim.y 419 #define hipGridDim_z gridDim.z 421 #include <hip/hcc_detail/math_functions.h> 423 #if __HIP_HCC_COMPAT_MODE__ 425 #pragma push_macro("__DEFINE_HCC_FUNC") 426 #define __DEFINE_HCC_FUNC(hc_fun,hip_var) \ 427 inline __device__ __attribute__((always_inline)) uint hc_get_##hc_fun(uint i) { \ 436 __DEFINE_HCC_FUNC(workitem_id, threadIdx)
437 __DEFINE_HCC_FUNC(group_id, blockIdx)
438 __DEFINE_HCC_FUNC(group_size, blockDim)
439 __DEFINE_HCC_FUNC(num_groups, gridDim)
440 #pragma pop_macro("__DEFINE_HCC_FUNC") 442 extern "C" __device__
__attribute__((
const))
size_t __ockl_get_global_id(uint);
444 hc_get_workitem_absolute_id(
int dim)
446 return (uint)__ockl_get_global_id(dim);
452 #pragma push_macro("__CUDA__") 454 #include <__clang_cuda_math_forward_declares.h> 455 #include <__clang_cuda_complex_builtins.h> 456 #include <cuda_wrappers/algorithm> 457 #include <cuda_wrappers/complex> 458 #include <cuda_wrappers/new> 460 #pragma pop_macro("__CUDA__") 463 uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
464 uint32_t localWorkSizeX, uint32_t localWorkSizeY,
465 uint32_t localWorkSizeZ,
size_t sharedMemBytes,
466 hipStream_t hStream,
void** kernelParams,
void** extra,
471 hipError_t hipHccModuleLaunchKernel(
hipFunction_t f, uint32_t globalWorkSizeX,
472 uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
473 uint32_t localWorkSizeX, uint32_t localWorkSizeY,
474 uint32_t localWorkSizeZ,
size_t sharedMemBytes,
475 hipStream_t hStream,
void** kernelParams,
void** extra,
478 __attribute__((deprecated(
"use hipExtModuleLaunchKernel instead")));
480 #endif // defined(__clang__) && defined(__HIP__) 482 #include <hip/hcc_detail/hip_memory.h> 484 #endif // HIP_HCC_DETAIL_RUNTIME_H
_Float16 __2f16 __attribute__((ext_vector_type(2)))
Copies the size of symbol symbolName to size.
Definition: hip_fp16_math_fwd.h:53
hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ, uint32_t localWorkSizeX, uint32_t localWorkSizeY, uint32_t localWorkSizeZ, size_t sharedMemBytes, hipStream_t hStream, void **kernelParams, void **extra, hipEvent_t startEvent=nullptr, hipEvent_t stopEvent=nullptr, uint32_t flags=0)
launches kernel f with launch parameters and shared memory on stream with arguments passed to kernelp...
Definition: hip_module.cpp:290
Definition: hip_module.cpp:106
Definition: hip_runtime_api.h:269
#define __host__
Definition: host_defines.h:41
Definition: grid_launch.h:31
Definition: hip_runtime_api.h:82
Definition: hip_hcc_internal.h:705
Definition: hip_hcc_internal.h:524
Definition: hip_runtime.h:204