char-fork/llama.cpp

git clone https://git.t4t.associates/char-fork/llama.cpp

AuroraRASggml: add gfx90c HIP support (#26454)473599738

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
11.9 KiB301 linesraw
1#pragma once
2
3#define HIP_DISABLE_WARP_SYNC_BUILTINS 1
4#include <hip/hip_runtime.h>
5#include <hipblas/hipblas.h>
6#include <hip/hip_fp16.h>
7#include <hip/hip_bf16.h>
8
9#ifdef GGML_USE_NCCL
10#include <rccl/rccl.h>
11#endif // GGML_USE_NCCL
12
13
14#define CUBLAS_GEMM_DEFAULT HIPBLAS_GEMM_DEFAULT
15#define CUBLAS_GEMM_DEFAULT_TENSOR_OP HIPBLAS_GEMM_DEFAULT
16#define CUBLAS_OP_N HIPBLAS_OP_N
17#define CUBLAS_OP_T HIPBLAS_OP_T
18#define CUBLAS_STATUS_SUCCESS HIPBLAS_STATUS_SUCCESS
19#define CUBLAS_TF32_TENSOR_OP_MATH 0
20#define CUDA_R_16F  HIPBLAS_R_16F
21#define CUDA_R_16BF HIPBLAS_R_16B
22#define CUDA_R_32F  HIPBLAS_R_32F
23#define CUBLAS_SIDE_RIGHT HIPBLAS_SIDE_RIGHT
24#define CUBLAS_FILL_MODE_UPPER HIPBLAS_FILL_MODE_UPPER
25#define CUBLAS_DIAG_NON_UNIT HIPBLAS_DIAG_NON_UNIT
26#define CU_DEVICE_ATTRIBUTE_VIRTUAL_MEMORY_MANAGEMENT_SUPPORTED hipDeviceAttributeVirtualMemoryManagementSupported
27#define CU_MEM_ALLOC_GRANULARITY_RECOMMENDED hipMemAllocationGranularityRecommended
28#define CU_MEM_ALLOCATION_TYPE_PINNED hipMemAllocationTypePinned
29#define CU_MEM_LOCATION_TYPE_DEVICE hipMemLocationTypeDevice
30#define CU_MEM_ACCESS_FLAGS_PROT_READWRITE hipMemAccessFlagsProtReadWrite
31#define CU_CHECK(fn) {hipError_t err = fn; if(err != hipSuccess) { GGML_ABORT("HipVMM Failure: %s\n", hipGetErrorString(err)); }}
32#define __shfl_sync(mask, var, laneMask, width) __shfl(var, laneMask, width)
33#define __shfl_up_sync(mask, var, laneMask, width) __shfl_up(var, laneMask, width)
34#define __shfl_xor_sync(mask, var, laneMask, width) __shfl_xor(var, laneMask, width)
35#define __all_sync(mask, var) __all(var)
36#define __any_sync(mask, var) __any(var)
37#define cublasStrsmBatched hipblasStrsmBatched
38#define cublasCreate hipblasCreate
39#define cublasDestroy hipblasDestroy
40#define cublasGemmEx hipblasGemmEx
41#define cublasGemmBatchedEx hipblasGemmBatchedEx
42#define cublasGemmStridedBatchedEx hipblasGemmStridedBatchedEx
43#define cublasHandle_t hipblasHandle_t
44#define cublasSetMathMode(handle, mode) CUBLAS_STATUS_SUCCESS
45#define cublasSetStream hipblasSetStream
46#define cublasSgemm hipblasSgemm
47#define cublasSgemmBatched hipblasSgemmBatched
48#define cublasSgemmStridedBatched hipblasSgemmStridedBatched
49#define cublasStatus_t hipblasStatus_t
50#define cublasOperation_t hipblasOperation_t
51#define cudaDevAttrCooperativeLaunch hipDeviceAttributeCooperativeLaunch
52#define cudaDeviceCanAccessPeer hipDeviceCanAccessPeer
53#define cudaDeviceDisablePeerAccess hipDeviceDisablePeerAccess
54#define cudaDeviceEnablePeerAccess hipDeviceEnablePeerAccess
55#define cudaDeviceGetAttribute hipDeviceGetAttribute
56#define cudaDeviceGetPCIBusId hipDeviceGetPCIBusId
57#define cudaDeviceProp hipDeviceProp_t
58#define cudaDeviceSynchronize hipDeviceSynchronize
59#define cudaError_t hipError_t
60#define cudaErrorMemoryAllocation hipErrorOutOfMemory
61#define cudaErrorPeerAccessAlreadyEnabled hipErrorPeerAccessAlreadyEnabled
62#define cudaErrorPeerAccessNotEnabled hipErrorPeerAccessNotEnabled
63#define cudaEventCreateWithFlags hipEventCreateWithFlags
64#define cudaEventDisableTiming hipEventDisableTiming
65#define cudaEventRecord hipEventRecord
66#define cudaEventSynchronize hipEventSynchronize
67#define cudaEvent_t hipEvent_t
68#define cudaEventDestroy hipEventDestroy
69#define cudaFree hipFree
70#define cudaFreeHost hipHostFree
71#define cudaGetDevice hipGetDevice
72#define cudaGetDeviceCount hipGetDeviceCount
73#define cudaGetDeviceProperties hipGetDeviceProperties
74#define cudaGetErrorString hipGetErrorString
75#define cudaGetLastError hipGetLastError
76#define cudaHostRegister hipHostRegister
77#define cudaHostRegisterPortable hipHostRegisterPortable
78#define cudaHostRegisterReadOnly hipHostRegisterReadOnly
79#define cudaHostUnregister hipHostUnregister
80#define cudaLaunchCooperativeKernel hipLaunchCooperativeKernel
81#define cudaLaunchHostFunc hipLaunchHostFunc
82#define cudaMalloc hipMalloc
83#define cudaMallocHost(ptr, size) hipHostMalloc(ptr, size, hipHostMallocDefault)
84#define cudaMallocManaged hipMallocManaged
85#define cudaMemAdvise hipMemAdvise
86#define cudaMemcpy hipMemcpy
87#define cudaMemcpyAsync hipMemcpyAsync
88#define cudaMemcpyPeerAsync hipMemcpyPeerAsync
89#define cudaMemcpy2DAsync hipMemcpy2DAsync
90#define cudaMemcpyDeviceToDevice hipMemcpyDeviceToDevice
91#define cudaMemcpyDeviceToHost hipMemcpyDeviceToHost
92#define cudaMemcpyHostToDevice hipMemcpyHostToDevice
93#define cudaMemcpyKind hipMemcpyKind
94#define cudaMemset hipMemset
95#define cudaMemsetAsync hipMemsetAsync
96#define cudaMemGetInfo hipMemGetInfo
97#define cudaOccupancyMaxPotentialBlockSize hipOccupancyMaxPotentialBlockSize
98#define cudaSetDevice hipSetDevice
99#define cuDeviceGet hipDeviceGet
100#define CUdevice hipDevice_t
101#define CUdeviceptr hipDeviceptr_t
102#define cuMemUnmap hipMemUnmap
103#define CUmemAccessDesc hipMemAccessDesc
104#define cuMemAddressFree hipMemAddressFree
105#define cuMemRelease hipMemRelease
106#define CUmemGenericAllocationHandle hipMemGenericAllocationHandle_t
107#define cuMemCreate hipMemCreate
108#define cuMemAddressReserve hipMemAddressReserve
109#define cuMemMap hipMemMap
110#define cuMemSetAccess hipMemSetAccess
111#define cuMemGetAllocationGranularity hipMemGetAllocationGranularity
112#define CUmemAllocationProp hipMemAllocationProp
113#define cuDeviceGetAttribute hipDeviceGetAttribute
114#define cudaStreamCreateWithFlags hipStreamCreateWithFlags
115#define cudaStreamDestroy hipStreamDestroy
116#define cudaStreamFireAndForget hipStreamFireAndForget
117#define cudaStreamNonBlocking hipStreamNonBlocking
118#define cudaStreamPerThread hipStreamPerThread
119#define cudaStreamSynchronize hipStreamSynchronize
120#define cudaStreamWaitEvent hipStreamWaitEvent
121#define cudaGraphExec_t hipGraphExec_t
122#define cudaGraphNode_t hipGraphNode_t
123#define cudaKernelNodeParams hipKernelNodeParams
124#define cudaKernelNodeParams hipKernelNodeParams
125#define cudaGraphExecDestroy hipGraphExecDestroy
126#define cudaGraphLaunch hipGraphLaunch
127#define cudaErrorGraphExecUpdateFailure hipErrorGraphExecUpdateFailure
128#define cudaGraphExecUpdateResult hipGraphExecUpdateResult
129#define cudaGraphNodeType hipGraphNodeType
130#define cudaGraphNodeTypeKernel hipGraphNodeTypeKernel
131#define cudaGraphInstantiate hipGraphInstantiate
132#define cudaStreamEndCapture hipStreamEndCapture
133#define cudaGraphDestroy hipGraphDestroy
134#define cudaGraphKernelNodeSetParams hipGraphKernelNodeSetParams
135#define cudaErrorInvalidDeviceFunction hipErrorInvalidDeviceFunction
136#define cudaGraphKernelNodeGetParams hipGraphKernelNodeGetParams
137#define cudaGraphNodeGetType hipGraphNodeGetType
138#define cudaGraphGetNodes hipGraphGetNodes
139#define cudaGraphExecUpdate hipGraphExecUpdate
140#define cudaStreamCaptureModeRelaxed hipStreamCaptureModeRelaxed
141#define cudaStreamBeginCapture hipStreamBeginCapture
142#define cudaGraph_t hipGraph_t
143#define cudaStream_t hipStream_t
144#define cudaSuccess hipSuccess
145#define cudaOccupancyMaxActiveBlocksPerMultiprocessor hipOccupancyMaxActiveBlocksPerMultiprocessor
146#define cudaFuncSetAttribute hipFuncSetAttribute
147#define cudaFuncAttributeMaxDynamicSharedMemorySize hipFuncAttributeMaxDynamicSharedMemorySize
148#define __trap() do { abort(); __builtin_unreachable(); } while(0)
149#define CUBLAS_STATUS_SUCCESS HIPBLAS_STATUS_SUCCESS
150#define CUBLAS_STATUS_NOT_INITIALIZED HIPBLAS_STATUS_NOT_INITIALIZED
151#define CUBLAS_STATUS_ALLOC_FAILED HIPBLAS_STATUS_ALLOC_FAILED
152#define CUBLAS_STATUS_INVALID_VALUE HIPBLAS_STATUS_INVALID_VALUE
153#define CUBLAS_STATUS_ARCH_MISMATCH HIPBLAS_STATUS_ARCH_MISMATCH
154#define CUBLAS_STATUS_MAPPING_ERROR HIPBLAS_STATUS_MAPPING_ERROR
155#define CUBLAS_STATUS_EXECUTION_FAILED HIPBLAS_STATUS_EXECUTION_FAILED
156#define CUBLAS_STATUS_INTERNAL_ERROR HIPBLAS_STATUS_INTERNAL_ERROR
157#define CUBLAS_STATUS_NOT_SUPPORTED HIPBLAS_STATUS_NOT_SUPPORTED
158
159#if HIP_VERSION >= 60500000
160#define CUBLAS_COMPUTE_16F HIPBLAS_COMPUTE_16F
161#define CUBLAS_COMPUTE_32F HIPBLAS_COMPUTE_32F
162#define CUBLAS_COMPUTE_32F_FAST_16F HIPBLAS_COMPUTE_32F_FAST_16F
163#define cublasComputeType_t hipblasComputeType_t
164#define cudaDataType_t hipDataType
165#else
166#define CUBLAS_COMPUTE_16F HIPBLAS_R_16F
167#define CUBLAS_COMPUTE_32F HIPBLAS_R_32F
168#define CUBLAS_COMPUTE_32F_FAST_16F HIPBLAS_R_32F
169#define cublasComputeType_t hipblasDatatype_t
170#define cudaDataType_t hipblasDatatype_t
171#endif // HIP_VERSION >= 6050000
172
173#if !defined(__HIP_PLATFORM_AMD__)
174#error "The HIP backend supports only AMD targets"
175#endif // !defined(__HIP_PLATFORM_AMD__)
176
177#define __CUDA_ARCH__ 1300
178
179#if defined(__gfx900__) || defined(__gfx906__) || defined(__gfx909__) || defined(__gfx90c__)
180#define GCN5
181#endif // defined(__gfx900__) || defined(__gfx906__) || defined(__gfx909__) || defined(__gfx90c__)
182
183#if defined(__gfx803__)
184#define GCN4
185#endif // defined(__gfx803__)
186
187#if defined(GCN5) || defined(GCN4)
188#define GCN
189#endif // defined(GCN5) || defined(GCN4)
190
191#if defined(__gfx950__)
192#define CDNA4
193#endif // defined(__gfx950__)
194
195#if defined(__gfx942__)
196#define CDNA3
197#endif // defined(__gfx942__)
198
199#if defined(__gfx90a__)
200#define CDNA2
201#endif // defined(__gfx90a__)
202
203#if defined(__gfx908__)
204#define CDNA1
205#endif // defined(__gfx908__)
206
207#if defined(CDNA4) || defined(CDNA3) || defined(CDNA2) || defined(CDNA1)
208#define CDNA // For the entire family
209#endif // defined(CDNA4) || defined(CDNA3) || defined(CDNA2) || defined(CDNA1)
210
211#if defined(__GFX12__)
212#define RDNA4
213#endif // defined(__GFX12__)
214
215#if defined(__GFX11__)
216#define RDNA3
217#endif // defined(__GFX11__)
218
219#if defined(__gfx1150__) || defined(__gfx1151__) || defined(__gfx1152__) || defined(__gfx1153__)
220#define RDNA3_5
221#endif // defined(__gfx1150__) || defined(__gfx1151__) || defined(__gfx1152__) || defined(__gfx1153__)
222
223#if defined(RDNA3) && !defined(RDNA3_5)
224#define RDNA3_0
225#endif // defined(RDNA3) && !defined(RDNA3_5)
226
227#if defined(__gfx1030__) || defined(__gfx1031__) || defined(__gfx1032__) || defined(__gfx1033__) || \
228    defined(__gfx1034__) || defined(__gfx1035__) || defined(__gfx1036__) || defined(__gfx1037__)
229#define RDNA2
230#endif
231
232#if defined(__gfx1010__) || defined(__gfx1012__)
233#define RDNA1
234#endif // defined(__gfx1010__) || defined(__gfx1012__)
235
236#if defined(RDNA4) || defined(RDNA3) || defined(RDNA2) || defined(RDNA1)
237#define RDNA // For the entire family
238#endif // defined(RDNA4) || defined(RDNA3) || defined(RDNA2) || defined(RDNA1)
239
240#ifndef __has_builtin
241    #define __has_builtin(x) 0
242#endif
243
244typedef __hip_bfloat16 nv_bfloat16;
245typedef __hip_bfloat162 nv_bfloat162;
246
247#if HIP_VERSION >= 60200000
248#include <hip/hip_fp8.h>
249typedef __hip_fp8_e4m3 __nv_fp8_e4m3;
250#define FP8_AVAILABLE
251#endif // HIP_VERSION >= 60200000
252
253typedef int8_t int8x4_t __attribute__((ext_vector_type(4)));
254typedef uint8_t uint8x4_t __attribute__((ext_vector_type(4)));
255static __device__ __forceinline__ int __vsubss4(const int a, const int b) {
256    const int8x4_t va = reinterpret_cast<const int8x4_t&>(a);
257    const int8x4_t vb = reinterpret_cast<const int8x4_t&>(b);
258#if __has_builtin(__builtin_elementwise_sub_sat)
259    const int8x4_t c = __builtin_elementwise_sub_sat(va, vb);
260    return reinterpret_cast<const int &>(c);
261#else
262    int8x4_t c;
263    int16_t tmp;
264#pragma unroll
265    for (int i = 0; i < 4; i++) {
266        tmp = va[i] - vb[i];
267        if(tmp > std::numeric_limits<int8_t>::max()) tmp = std::numeric_limits<int8_t>::max();
268        if(tmp < std::numeric_limits<int8_t>::min()) tmp = std::numeric_limits<int8_t>::min();
269        c[i] = tmp;
270    }
271    return reinterpret_cast<int &>(c);
272#endif // __has_builtin(__builtin_elementwise_sub_sat)
273}
274
275static __device__ __forceinline__ int __vsub4(const int a, const int b) {
276    return __vsubss4(a, b);
277}
278
279static __device__ __forceinline__ unsigned int __vcmpeq4(unsigned int a, unsigned int b) {
280    const uint8x4_t& va = reinterpret_cast<const uint8x4_t&>(a);
281    const uint8x4_t& vb = reinterpret_cast<const uint8x4_t&>(b);
282    unsigned int c;
283    uint8x4_t& vc = reinterpret_cast<uint8x4_t&>(c);
284#pragma unroll
285    for (int i = 0; i < 4; ++i) {
286        vc[i] = va[i] == vb[i] ? 0xff : 0x00;
287    }
288    return c;
289}
290
291static __device__ __forceinline__ unsigned int __vcmpne4(unsigned int a, unsigned int b) {
292    const uint8x4_t& va = reinterpret_cast<const uint8x4_t&>(a);
293    const uint8x4_t& vb = reinterpret_cast<const uint8x4_t&>(b);
294    unsigned int c;
295    uint8x4_t& vc = reinterpret_cast<uint8x4_t&>(c);
296#pragma unroll
297    for (int i = 0; i < 4; ++i) {
298        vc[i] = va[i] == vb[i] ? 0x00 : 0xff;
299    }
300    return c;
301}