char-fork/llama.cpp

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

Johannes GäßlerCUDA: fix BF16 FA compilation (#20865)bd3f1d9d6

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
2.3 KiB66 linesraw
1#pragma once
2#include "common.cuh"
3
4#define CUDA_DEQUANTIZE_BLOCK_SIZE 256
5
6template<typename T>
7using to_t_cuda_t = void (*)(const void * x, T * y, int64_t k, cudaStream_t stream);
8
9typedef to_t_cuda_t<float> to_fp32_cuda_t;
10typedef to_t_cuda_t<half> to_fp16_cuda_t;
11typedef to_t_cuda_t<nv_bfloat16> to_bf16_cuda_t;
12
13to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type);
14
15to_bf16_cuda_t ggml_get_to_bf16_cuda(ggml_type type);
16
17to_fp32_cuda_t ggml_get_to_fp32_cuda(ggml_type type);
18
19// TODO more general support for non-contiguous inputs
20
21template<typename T>
22using to_t_nc_cuda_t = void (*)(const void * x, T * y,
23    int64_t ne00, int64_t ne01, int64_t ne02, int64_t ne03,
24    int64_t s01, int64_t s02, int64_t s03, cudaStream_t stream);
25
26typedef to_t_nc_cuda_t<float> to_fp32_nc_cuda_t;
27typedef to_t_nc_cuda_t<half> to_fp16_nc_cuda_t;
28typedef to_t_nc_cuda_t<nv_bfloat16> to_bf16_nc_cuda_t;
29
30to_fp32_nc_cuda_t ggml_get_to_fp32_nc_cuda(ggml_type type);
31to_fp16_nc_cuda_t ggml_get_to_fp16_nc_cuda(ggml_type type);
32to_bf16_nc_cuda_t ggml_get_to_bf16_nc_cuda(ggml_type type);
33
34template<typename dst_t, typename src_t>
35 __host__ __device__ inline dst_t ggml_cuda_cast(src_t x) {
36    if constexpr (std::is_same_v<dst_t, src_t>) {
37        return x;
38    } else if constexpr(std::is_same_v<dst_t, nv_bfloat16>) {
39        return __float2bfloat16(float(x));
40    } else if constexpr(std::is_same_v<src_t, nv_bfloat16>) {
41        return __bfloat162float(x);
42    } else if constexpr(std::is_same_v<src_t, float2> && std::is_same_v<dst_t, half2>) {
43        return __float22half2_rn(x);
44    } else if constexpr(std::is_same_v<src_t, nv_bfloat162> && std::is_same_v<dst_t, float2>) {
45#ifdef GGML_USE_HIP
46        return make_float2(__bfloat162float(__low2bfloat16(x)), __bfloat162float(__high2bfloat16(x)));
47#else
48#if __CUDA_ARCH__ >= 800
49        return __bfloat1622float2(x);
50#else
51        return make_float2(__bfloat162float(x.x), __bfloat162float(x.y));
52#endif // __CUDA_ARCH__ >= 800
53#endif // GGML_USE_HIP
54    } else if constexpr(std::is_same_v<src_t, float2> && std::is_same_v<dst_t, nv_bfloat162>) {
55        // bypass compile error on cuda 12.0.1
56#ifdef GGML_USE_HIP
57        return __float22bfloat162_rn(x);
58#else
59        return {x.x, x.y};
60#endif // GGML_USE_HIP
61    } else if constexpr(std::is_same_v<dst_t, int32_t>) {
62        return int32_t(x);
63    } else {
64        return float(x);
65    }
66}