mirror of
https://github.com/huggingface/candle.git
synced 2025-06-15 18:28:24 +00:00
Merge pull request #601 from huggingface/repair_bf16_f16_cast
Repairing cast bf16/f16
This commit is contained in:
@ -24,19 +24,42 @@ extern "C" __global__ void FN_NAME( \
|
||||
} \
|
||||
} \
|
||||
|
||||
#define CAST_BF_OP(SRC_TYPENAME, DST_TYPENAME, FN_NAME) \
|
||||
extern "C" __global__ void FN_NAME( \
|
||||
const size_t numel, \
|
||||
const size_t num_dims, \
|
||||
const size_t *info, \
|
||||
const SRC_TYPENAME *inp, \
|
||||
DST_TYPENAME *out \
|
||||
) { \
|
||||
const size_t *dims = info; \
|
||||
const size_t *strides = info + num_dims; \
|
||||
if (is_contiguous(num_dims, dims, strides)) { \
|
||||
for (unsigned int i = blockIdx.x * blockDim.x + threadIdx.x; i < numel; i += blockDim.x * gridDim.x) { \
|
||||
out[i] = (DST_TYPENAME) (float) inp[i]; \
|
||||
} \
|
||||
} \
|
||||
else { \
|
||||
for (unsigned int i = blockIdx.x * blockDim.x + threadIdx.x; i < numel; i += blockDim.x * gridDim.x) { \
|
||||
unsigned strided_i = get_strided_index(i, num_dims, dims, strides); \
|
||||
out[i] = (DST_TYPENAME) (float) inp[strided_i]; \
|
||||
} \
|
||||
} \
|
||||
} \
|
||||
|
||||
#if __CUDA_ARCH__ >= 800
|
||||
CAST_OP(__nv_bfloat16, __nv_bfloat16, cast_bf16_bf16)
|
||||
|
||||
// CAST_OP(__nv_bfloat16, uint8_t, cast_bf16_u8)
|
||||
CAST_OP(__nv_bfloat16, uint32_t, cast_bf16_u32)
|
||||
// CAST_OP(__nv_bfloat16, __half, cast_bf16_f16)
|
||||
CAST_OP(__nv_bfloat16, float, cast_bf16_f32)
|
||||
CAST_OP(__nv_bfloat16, double, cast_bf16_f64)
|
||||
CAST_OP(uint8_t, __nv_bfloat16, cast_u8_bf16)
|
||||
CAST_OP(uint32_t, __nv_bfloat16, cast_u32_bf16)
|
||||
// CAST_OP(__half, __nv_bfloat16, cast_f16_bf16)
|
||||
CAST_OP(float, __nv_bfloat16, cast_f32_bf16)
|
||||
CAST_OP(double, __nv_bfloat16, cast_f64_bf16)
|
||||
CAST_BF_OP(__nv_bfloat16, __half, cast_bf16_f16)
|
||||
CAST_BF_OP(__half, __nv_bfloat16, cast_f16_bf16)
|
||||
#endif
|
||||
|
||||
#if __CUDA_ARCH__ >= 530
|
||||
|
Reference in New Issue
Block a user