HIP: Heterogenous-computing Interface for Portability
hip_runtime.h
1 /*
2 Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
3 
4 Permission is hereby granted, free of charge, to any person obtaining a copy
5 of this software and associated documentation files (the "Software"), to deal
6 in the Software without restriction, including without limitation the rights
7 to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
8 copies of the Software, and to permit persons to whom the Software is
9 furnished to do so, subject to the following conditions:
10 
11 The above copyright notice and this permission notice shall be included in
12 all copies or substantial portions of the Software.
13 
14 THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
15 IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
16 FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
17 AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
18 LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
19 OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
20 THE SOFTWARE.
21 */
22 
23 #ifndef HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_RUNTIME_H
24 #define HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_RUNTIME_H
25 
26 #include <cuda_runtime.h>
27 
28 #include <hip/hip_runtime_api.h>
29 
30 #define HIP_KERNEL_NAME(...) __VA_ARGS__
31 
32 typedef int hipLaunchParm;
33 
34 #define hipLaunchKernelGGLInternal(kernelName, numBlocks, numThreads, memPerBlock, streamId, ...) \
35  do { \
36  kernelName<<<numBlocks, numThreads, memPerBlock, streamId>>>(__VA_ARGS__); \
37  } while (0)
38 
39 #define hipLaunchKernelGGL(kernelName, ...) hipLaunchKernelGGLInternal((kernelName), __VA_ARGS__)
40 
41 #define hipReadModeElementType cudaReadModeElementType
42 
43 #ifdef __CUDA_ARCH__
44 
45 
46 // 32-bit Atomics:
47 #define __HIP_ARCH_HAS_GLOBAL_INT32_ATOMICS__ (__CUDA_ARCH__ >= 110)
48 #define __HIP_ARCH_HAS_GLOBAL_FLOAT_ATOMIC_EXCH__ (__CUDA_ARCH__ >= 110)
49 #define __HIP_ARCH_HAS_SHARED_INT32_ATOMICS__ (__CUDA_ARCH__ >= 120)
50 #define __HIP_ARCH_HAS_SHARED_FLOAT_ATOMIC_EXCH__ (__CUDA_ARCH__ >= 120)
51 #define __HIP_ARCH_HAS_FLOAT_ATOMIC_ADD__ (__CUDA_ARCH__ >= 200)
52 
53 // 64-bit Atomics:
54 #define __HIP_ARCH_HAS_GLOBAL_INT64_ATOMICS__ (__CUDA_ARCH__ >= 200)
55 #define __HIP_ARCH_HAS_SHARED_INT64_ATOMICS__ (__CUDA_ARCH__ >= 120)
56 
57 // Doubles
58 #define __HIP_ARCH_HAS_DOUBLES__ (__CUDA_ARCH__ >= 120)
59 
60 // warp cross-lane operations:
61 #define __HIP_ARCH_HAS_WARP_VOTE__ (__CUDA_ARCH__ >= 120)
62 #define __HIP_ARCH_HAS_WARP_BALLOT__ (__CUDA_ARCH__ >= 200)
63 #define __HIP_ARCH_HAS_WARP_SHUFFLE__ (__CUDA_ARCH__ >= 300)
64 #define __HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__ (__CUDA_ARCH__ >= 350)
65 
66 // sync
67 #define __HIP_ARCH_HAS_THREAD_FENCE_SYSTEM__ (__CUDA_ARCH__ >= 200)
68 #define __HIP_ARCH_HAS_SYNC_THREAD_EXT__ (__CUDA_ARCH__ >= 200)
69 
70 // misc
71 #define __HIP_ARCH_HAS_SURFACE_FUNCS__ (__CUDA_ARCH__ >= 200)
72 #define __HIP_ARCH_HAS_3DGRID__ (__CUDA_ARCH__ >= 200)
73 #define __HIP_ARCH_HAS_DYNAMIC_PARALLEL__ (__CUDA_ARCH__ >= 350)
74 
75 #endif
76 
77 #ifdef __CUDACC__
78 
79 
80 #define hipThreadIdx_x threadIdx.x
81 #define hipThreadIdx_y threadIdx.y
82 #define hipThreadIdx_z threadIdx.z
83 
84 #define hipBlockIdx_x blockIdx.x
85 #define hipBlockIdx_y blockIdx.y
86 #define hipBlockIdx_z blockIdx.z
87 
88 #define hipBlockDim_x blockDim.x
89 #define hipBlockDim_y blockDim.y
90 #define hipBlockDim_z blockDim.z
91 
92 #define hipGridDim_x gridDim.x
93 #define hipGridDim_y gridDim.y
94 #define hipGridDim_z gridDim.z
95 
96 #define HIP_SYMBOL(X) &X
97 
102 #define HIP_DYNAMIC_SHARED(type, var) extern __shared__ type var[];
103 
104 #define HIP_DYNAMIC_SHARED_ATTRIBUTE
105 
106 #ifdef __HIP_DEVICE_COMPILE__
107 #define abort_() \
108  { asm("trap;"); }
109 #undef assert
110 #define assert(COND) \
111  { \
112  if (!COND) { \
113  abort_(); \
114  } \
115  }
116 #endif
117 
118 #define __clock() clock()
119 #define __clock64() clock64()
120 
121 #endif
122 
123 #endif
hip_runtime_api.h
Defines the API signatures for HIP runtime. This file can be compiled with a standard compiler.