more fp8

Author: zackangelo

candle-core/benches/benchmarks/affine.rs

@@ -37,6 +37,7 @@ fn criterion_benchmark(c: &mut Criterion) {
         run_affine_benchmark(c, &device, DType::F32, "affine_f32");
         run_affine_benchmark(c, &device, DType::F16, "affine_f16");
         run_affine_benchmark(c, &device, DType::BF16, "affine_bf16");
+        run_affine_benchmark(c, &device, DType::F8E4M3, "affine_fp8");
     }
 }
 

candle-core/benches/benchmarks/mod.rs

@@ -8,7 +8,7 @@ pub(crate) mod reduce;
 pub(crate) mod unary;
 pub(crate) mod where_cond;
 
-use candle_core::{Device, Result};
+use candle_core::{backend::BackendDevice, Device, Result};
 
 pub(crate) trait BenchDevice {
     fn sync(&self) -> Result<()>;

candle-kernels/src/affine.cu

@@ -30,7 +30,9 @@ extern "C" __global__ void FN_NAME(  \
 
 #if __CUDA_ARCH__ >= 800
 AFFINE_OP(__nv_bfloat16, affine_bf16, x * mul + add)
+#endif
 
+#if __CUDA_ARCH__ >= 890
 #define F8E4M3_TO_FLOAT(x) __half2float(__nv_cvt_fp8_to_halfraw(x.__x, __NV_E4M3))
 
 AFFINE_OP(__nv_fp8_e4m3, affine_f8_e4m3, __nv_fp8_e4m3(F8E4M3_TO_FLOAT(x) * F8E4M3_TO_FLOAT(mul) + F8E4M3_TO_FLOAT(add)))