huggingface/candle
1
0
Fork 0
mirror of https://github.com/huggingface/candle.git synced 2025-06-17 19:18:50 +00:00
Code Issues Packages Projects Releases Wiki Activity

Intermediary float cast is necessary for cuda 11.8

Browse Source
This commit is contained in:
Nicolas Patry
2023-08-25 11:54:30 +00:00
parent 1c1e34735e
commit be371e827c
1 changed files with 2 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
candle-kernels/src/cast.cu
View File

@ -13,13 +13,13 @@ extern "C" __global__ void FN_NAME( \
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] = static_cast<DST_TYPENAME>(inp[i]); \
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] = static_cast<DST_TYPENAME>(inp[strided_i]); \
out[i] = (DST_TYPENAME) (float) inp[strided_i]; \
} \
} \
} \