Skip to content

Commit 0267ef1

Browse files
haricotharicot
committed
resolved conflicts and fp8 fallback for 530--870
1 parent 91dca22 commit 0267ef1

6 files changed

Lines changed: 80 additions & 55 deletions

File tree

‎candle-core/tests/tensor_tests.rs‎

Lines changed: 46 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1946,3 +1946,49 @@ fn tensor_norm() -> Result<()> {
19461946
assert_eq!(norm.to_scalar::<f64>()?, 5.);
19471947
Ok(())
19481948
}
1949+
1950+
fn affine_bf16(device: &Device) -> Result<()> {
1951+
// Test BF16 (native or fallback if cuda_arch < 700)
1952+
let bf16_vals = vec![
1953+
half::bf16::from_f32(1.0),
1954+
half::bf16::from_f32(2.0),
1955+
half::bf16::from_f32(3.0),
1956+
];
1957+
let tensor_bf16 = Tensor::new(bf16_vals, device)?;
1958+
let affine_bf16 = tensor_bf16.affine(2.0, 1.0)?;
1959+
let expected_bf16 = vec![
1960+
half::bf16::from_f32(3.0),
1961+
half::bf16::from_f32(5.0),
1962+
half::bf16::from_f32(7.0),
1963+
];
1964+
let expected_tensor_bf16 = Tensor::new(expected_bf16, device)?;
1965+
candle_core::test_utils::assert_tensor_eq(&affine_bf16, &expected_tensor_bf16)?;
1966+
Ok(())
1967+
}
1968+
1969+
fn affine_fp8(device: &Device) -> Result<()> {
1970+
// Test F8E4M3 (native or fallback if cuda_arch <= 890)
1971+
let f8_vals = vec![
1972+
F8E4M3::from_f32(1.0),
1973+
F8E4M3::from_f32(2.0),
1974+
F8E4M3::from_f32(3.0),
1975+
];
1976+
let tensor_f8 = Tensor::new(f8_vals, device)?;
1977+
let affine_f8 = tensor_f8.affine(2.0, 1.0)?;
1978+
let expected_f8 = vec![
1979+
F8E4M3::from_f32(3.0),
1980+
F8E4M3::from_f32(5.0),
1981+
F8E4M3::from_f32(7.0),
1982+
];
1983+
let expected_tensor_f8 = Tensor::new(expected_f8, device)?;
1984+
candle_core::test_utils::assert_tensor_eq(&affine_f8, &expected_tensor_f8)?;
1985+
Ok(())
1986+
}
1987+
1988+
test_device!(
1989+
affine_bf16,
1990+
affine_bf16_cpu,
1991+
affine_bf16_gpu,
1992+
affine_bf16_metal
1993+
);
1994+
test_device!(affine_fp8, affine_fp8_cpu, affine_fp8_gpu, affine_fp8_metal);

‎candle-kernels/src/affine.cu‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -32,12 +32,11 @@ extern "C" __global__ void FN_NAME( \
3232
AFFINE_OP(__nv_bfloat16, affine_bf16, x * mul + add)
3333
#endif
3434

35-
#if __CUDA_ARCH__ >= 890
35+
#if __CUDA_ARCH__ >= 890 || (__CUDA_ARCH__ >= 530 && __CUDA_ARCH__ < 890)
3636
#define F8E4M3_TO_FLOAT(x) __half2float(__nv_cvt_fp8_to_halfraw(x.__x, __NV_E4M3))
3737

3838
AFFINE_OP(__nv_fp8_e4m3, affine_f8_e4m3, __nv_fp8_e4m3(F8E4M3_TO_FLOAT(x) * F8E4M3_TO_FLOAT(mul) + F8E4M3_TO_FLOAT(add)))
3939
#endif
40-
#endif
4140

4241
#if __CUDA_ARCH__ >= 530
4342
AFFINE_OP(__half, affine_f16, x * mul + add)

‎candle-kernels/src/cast.cu‎

Lines changed: 3 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -169,23 +169,11 @@ CAST_OP_FP8_INTO(int32_t, __nv_fp8_e4m3, cast_i32_f8_e4m3)
169169
CAST_OP_FP8(__nv_fp8_e4m3, int32_t, cast_f8_e4m3_i32)
170170
CAST_OP_FP8(__nv_fp8_e4m3, __nv_bfloat16, cast_f8_e4m3_bf16)
171171
CAST_OP_FP8_INTO(__nv_bfloat16, __nv_fp8_e4m3, cast_bf16_f8_e4m3)
172-
#else
173-
#include <cuda.h>
174-
#if CUDA_VERSION >= 11000
175-
CAST_OP(__nv_bfloat16, float, cast_bf16_f32)
176-
CAST_OP(float, __nv_bfloat16, cast_f32_bf16)
177-
CAST_THROUGH_OP(__nv_bfloat16, uint8_t, float, cast_bf16_u8)
178-
CAST_THROUGH_OP(__nv_bfloat16, __half, float, cast_bf16_f16)
179-
CAST_THROUGH_OP(__nv_bfloat16, double, float, cast_bf16_f64)
180-
CAST_THROUGH_OP(__half, __nv_bfloat16, float, cast_f16_bf16)
181-
CAST_THROUGH_OP(double, __nv_bfloat16, float, cast_f64_bf16)
182-
CAST_THROUGH_OP(uint8_t, __nv_bfloat16, float, cast_u8_bf16)
183-
CAST_THROUGH_OP(__nv_bfloat16, __nv_fp8_e4m3, float, cast_bf16_f8_e4m3)
184-
#endif
185-
#endif
186172
#endif
173+
174+
187175
#if __CUDA_ARCH__ >= 530
188-
CAST_OP(__half, __half, cast_f16_f16)
176+
CAST_OP(__half, __half, cast_f16_f16)
189177

190178
CAST_THROUGH_OP(__half, uint8_t, float, cast_f16_u8)
191179
CAST_OP(__half, uint32_t, cast_f16_u32)

‎candle-kernels/src/compatibility.cuh‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -35,7 +35,7 @@ __device__ double atomicAdd(double* address, double val) {
3535
}
3636
#endif
3737

38-
#if __CUDA_ARCH__ < 700
38+
#if __CUDA_ARCH__ >= 530 && __CUDA_ARCH__ < 700
3939
// https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicadd
4040
// The 16-bit __half floating-point version of atomicAdd() is only supported by devices of compute capability 7.x and higher.
4141
// Solution adapted from https://github.com/torch/cutorch/blob/master/lib/THC/THCAtomics.cuh#L96-L119

‎candle-kernels/src/cuda_utils.cuh‎

Lines changed: 21 additions & 23 deletions
Original file line numberDiff line numberDiff line change
@@ -159,6 +159,27 @@ __device__ __forceinline__ uint32_t maxg(uint32_t a, uint32_t b) { return max(a,
159159
__device__ __forceinline__ uint8_t ming(uint8_t a, uint8_t b) { return min(a, b); }
160160
__device__ __forceinline__ uint8_t maxg(uint8_t a, uint8_t b) { return max(a, b); }
161161

162+
#if __CUDA_ARCH__ >= 530
163+
__device__ __forceinline__ __half powg(__half a, __half b) { return __float2half(powf(__half2float(a), __half2float(b))); }
164+
__device__ __forceinline__ bool isnang(__half a) { return __hisnan(a); }
165+
__device__ __forceinline__ __half sqrtg(__half a) { return hsqrt(a); }
166+
__device__ __forceinline__ __half cosg(__half a) { return hcos(a); }
167+
__device__ __forceinline__ __half sing(__half a) { return hsin(a); }
168+
__device__ __forceinline__ __half recipg(__half a) { __half one = 1.0; return one / a; }
169+
__device__ __forceinline__ __half maxg(__half a, __half b) { return __hmax_nan(a, b); }
170+
__device__ __forceinline__ __half tanhg(__half a) { return __float2half(tanhf(__half2float(a))); }
171+
__device__ __forceinline__ __half erfg(__half a) { return __float2half(erff(__half2float(a))); }
172+
__device__ __forceinline__ __half ceilg(__half a) { return __float2half(ceilf(__half2float(a))); }
173+
__device__ __forceinline__ __half floorg(__half a) { return __float2half(floorf(__half2float(a))); }
174+
__device__ __forceinline__ __half roundg(__half a) { return __float2half(roundf(__half2float(a))); }
175+
__device__ __forceinline__ __half normcdfg(__half a) { return __float2half(normcdff(__half2float(a))); }
176+
__device__ __forceinline__ __half ming(__half a, __half b) { return __hmin_nan(a, b); }
177+
__device__ __forceinline__ __half logg(__half a) { return hlog(a); }
178+
__device__ __forceinline__ __half expg(__half a) { return hexp(a); }
179+
__device__ __forceinline__ __half absg(__half a) { return __habs(a); }
180+
__device__ __forceinline__ __half copysigng(__half a, __half b) { return __float2half(copysignf(__half2float(a), __half2float(b))); }
181+
#endif
182+
162183
#if __CUDA_ARCH__ >= 800 || (__CUDA_ARCH__ >= 530 && __CUDA_ARCH__ < 800)
163184
__device__ __forceinline__ __nv_bfloat16 powg(__nv_bfloat16 a, __nv_bfloat16 b) { return __float2bfloat16(powf(__bfloat162float(a), __bfloat162float(b))); }
164185
__device__ __forceinline__ bool isnang(__nv_bfloat16 a) { return __hisnan(a); }
@@ -199,27 +220,4 @@ __device__ __forceinline__ __nv_fp8_e4m3 logg(__nv_fp8_e4m3 a) { return __nv_fp8
199220
__device__ __forceinline__ __nv_fp8_e4m3 expg(__nv_fp8_e4m3 a) { return __nv_fp8_e4m3(expf(F8E4M3_TO_FLOAT(a))); }
200221
__device__ __forceinline__ __nv_fp8_e4m3 absg(__nv_fp8_e4m3 a) { return __nv_fp8_e4m3(fabsf(F8E4M3_TO_FLOAT(a))); }
201222
__device__ __forceinline__ __nv_fp8_e4m3 copysigng(__nv_fp8_e4m3 a, __nv_fp8_e4m3 b) { return __nv_fp8_e4m3(copysignf(F8E4M3_TO_FLOAT(a), F8E4M3_TO_FLOAT(b))); }
202-
203-
204-
#endif
205-
206-
#if __CUDA_ARCH__ >= 530
207-
__device__ __forceinline__ __half powg(__half a, __half b) { return __float2half(powf(__half2float(a), __half2float(b))); }
208-
__device__ __forceinline__ bool isnang(__half a) { return __hisnan(a); }
209-
__device__ __forceinline__ __half sqrtg(__half a) { return hsqrt(a); }
210-
__device__ __forceinline__ __half cosg(__half a) { return hcos(a); }
211-
__device__ __forceinline__ __half sing(__half a) { return hsin(a); }
212-
__device__ __forceinline__ __half recipg(__half a) { __half one = 1.0; return one / a; }
213-
__device__ __forceinline__ __half maxg(__half a, __half b) { return __hmax_nan(a, b); }
214-
__device__ __forceinline__ __half tanhg(__half a) { return __float2half(tanhf(__half2float(a))); }
215-
__device__ __forceinline__ __half erfg(__half a) { return __float2half(erff(__half2float(a))); }
216-
__device__ __forceinline__ __half ceilg(__half a) { return __float2half(ceilf(__half2float(a))); }
217-
__device__ __forceinline__ __half floorg(__half a) { return __float2half(floorf(__half2float(a))); }
218-
__device__ __forceinline__ __half roundg(__half a) { return __float2half(roundf(__half2float(a))); }
219-
__device__ __forceinline__ __half normcdfg(__half a) { return __float2half(normcdff(__half2float(a))); }
220-
__device__ __forceinline__ __half ming(__half a, __half b) { return __hmin_nan(a, b); }
221-
__device__ __forceinline__ __half logg(__half a) { return hlog(a); }
222-
__device__ __forceinline__ __half expg(__half a) { return hexp(a); }
223-
__device__ __forceinline__ __half absg(__half a) { return __habs(a); }
224-
__device__ __forceinline__ __half copysigng(__half a, __half b) { return __float2half(copysignf(__half2float(a), __half2float(b))); }
225223
#endif

‎candle-nn/tests/ops.rs‎

Lines changed: 8 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -4,7 +4,7 @@ extern crate intel_mkl_src;
44
#[cfg(feature = "accelerate")]
55
extern crate accelerate_src;
66

7-
use candle::{test_device, test_utils::to_vec3_round, Device, DType, IndexOp, Result, Tensor};
7+
use candle::{test_device, test_utils::to_vec3_round, DType, Device, IndexOp, Result, Tensor};
88

99
fn softmax(device: &Device) -> Result<()> {
1010
let data = &[[[3f32, 1., 4.], [1., 5., 9.]], [[2., 1., 7.], [8., 2., 8.]]];
@@ -324,17 +324,8 @@ fn sigmoid(device: &Device) -> Result<()> {
324324
Ok(())
325325
}
326326

327-
fn sigmoid_f16(device: &Device) -> Result<()> {
328-
let data = &[[[3f32, 1., 4.], [1., 5., 9.]], [[2., 1., 7.], [8., 2., 8.]]];
329-
let tensor = Tensor::new(data, device)?.to_dtype(DType::F16)?;
330-
let s1 = candle_nn::ops::sigmoid(&tensor)?;
331-
let s2 = (1. / (1. + tensor.neg()?.exp()?)?)?;
332-
let diff = (s1 - s2)?.abs()?.sum_all()?.to_vec0::<half::f16>()?;
333-
assert_eq!(diff, half::f16::from_f32(0.));
334-
Ok(())
335-
}
336-
337327
fn sigmoid_bf16(device: &Device) -> Result<()> {
328+
// Test BF16 (native or fallback if cuda_arch < 700)
338329
let data = &[[[3f32, 1., 4.], [1., 5., 9.]], [[2., 1., 7.], [8., 2., 8.]]];
339330
let tensor = Tensor::new(data, device)?.to_dtype(DType::BF16)?;
340331
let s1 = candle_nn::ops::sigmoid(&tensor)?;
@@ -343,7 +334,6 @@ fn sigmoid_bf16(device: &Device) -> Result<()> {
343334
assert_eq!(diff, half::bf16::from_f32(0.));
344335
Ok(())
345336
}
346-
347337

348338
test_device!(ropei, ropei_cpu, ropei_gpu, ropei_metal);
349339
test_device!(rope, rope_cpu, rope_gpu, rope_metal);
@@ -354,5 +344,9 @@ test_device!(rms_norml, rms_norml_cpu, rms_norml_gpu, rms_norml_metal);
354344
test_device!(layer_norm, ln_cpu, ln_gpu, ln_metal);
355345
test_device!(layer_norml, lnl_cpu, lnl_gpu, lnl_metal);
356346
test_device!(sigmoid, sigmoid_cpu, sigmoid_gpu, sigmoid_metal);
357-
test_device!(sigmoid_f16, sigmoid_b16_cpu, sigmoid_b16_gpu, sigmoid_b16_metal);
358-
test_device!(sigmoid_bf16, sigmoid_bf16_cpu, sigmoid_bf16_gpu, sigmoid_bf16_metal);
347+
test_device!(
348+
sigmoid_bf16,
349+
sigmoid_bf16_cpu,
350+
sigmoid_bf16_gpu,
351+
sigmoid_bf16_metal
352+
);

0 commit comments

Comments
 (0)