From 81459a1cffe1b1c67388e65b50a70bdd1c0926d9 Mon Sep 17 00:00:00 2001 From: Romaric Jodin Date: Wed, 1 Jul 2026 16:38:40 +0200 Subject: [PATCH] Add options to control denormal modes Introduce `-denorm-preserve` and `-denorm-flush-to-zero` options to clspv. These options allow specifying which floating-point types (fp16, fp32, fp64) should preserve denormals or flush them to zero. Refactor existing options to use a unified `FloatingPointType` enum instead of separate enums for FMA and RoundingModeRTE. Introduce `DenormMode` enum to represent the denormal handling state (preserve, flush_to_zero, unspecified, error). The new options are used to: - Set Clang's `FP32DenormalMode` and `FPDenormalMode` accordingly. - Emit `DenormPreserve` or `DenormFlushToZero` execution modes and capabilities in the SPIR-V producer. Validation is added in `ParseOptions` to ensure: - Preserve and FlushToZero are mutually exclusive for a given type. - fp16 and fp64 share the same mode, as they are mapped to the single `FPDenormalMode` in Clang CodeGenOptions. The unused `cl-denorms-are-zero` option is marked as having no effect. --- include/clspv/Option.h | 19 +++++---- lib/Compiler.cpp | 82 ++++++++++++++++++++++++++++++++++++--- lib/Option.cpp | 79 ++++++++++++++++++++++++++++++------- lib/SPIRVProducerPass.cpp | 72 ++++++++++++++++++++++++++++++---- test/denorm_modes.cl | 49 +++++++++++++++++++++++ 5 files changed, 269 insertions(+), 32 deletions(-) create mode 100644 test/denorm_modes.cl diff --git a/include/clspv/Option.h b/include/clspv/Option.h index 2faed9c0e..482bd6036 100644 --- a/include/clspv/Option.h +++ b/include/clspv/Option.h @@ -189,7 +189,7 @@ inline bool LanguageUsesGenericAddressSpace() { clspv::FeatureMacro::__opencl_c_generic_address_space) > 0); } -enum class RoundingModeRTE : uint32_t { +enum class FloatingPointType : uint32_t { fp16, fp32, fp64, @@ -197,7 +197,17 @@ enum class RoundingModeRTE : uint32_t { // Returns true when the execution mode RoundingModeRTE should be set for a // floating point type. -bool ExecutionModeRoundingModeRTE(RoundingModeRTE rm); +bool ExecutionModeRoundingModeRTE(FloatingPointType rm); + +enum class DenormMode : uint32_t { + error, + preserve, + flush_to_zero, + unspecified, +}; + +// Returns the execution mode for a floating point type. +DenormMode ExecutionModeDenorm(FloatingPointType fpty); // Return the SPIR-V binary version enum class SPIRVVersion : uint32_t { @@ -319,11 +329,6 @@ bool UntypedPointers(); // Returns true if untyped pointers are supported in aspace. bool UntypedPointerAddressSpace(unsigned aspace); -enum class SpvKhrFma : uint32_t { - fp16, - fp32, - fp64, -}; // Returns true if OpFmaKHR is supported for the type size. bool SupportsFmaKHR(uint32_t scalarSizeInBits); diff --git a/lib/Compiler.cpp b/lib/Compiler.cpp index 9ecbf30ab..e059dfd1e 100644 --- a/lib/Compiler.cpp +++ b/lib/Compiler.cpp @@ -93,6 +93,8 @@ static llvm::cl::opt cl_single_precision_constants( llvm::cl::desc("Treat double precision floating-point constant as single " "precision constant.")); +// This option has not effect. Mainly because of lack of expresiveness. Instead +// use `-denorm-preserve` and `-denorm-flush-to-zero`. static llvm::cl::opt cl_denorms_are_zero( "cl-denorms-are-zero", llvm::cl::init(false), llvm::cl::desc("If specified, denormalized floating point numbers may be " @@ -317,8 +319,52 @@ int SetCompilerInstanceOptions( instance.getLangOpts().NativeHalfType = true; instance.getLangOpts().NativeHalfArgsAndReturns = true; } - instance.getCodeGenOpts().FPDenormalMode = llvm::DenormalMode::getDynamic(); - instance.getCodeGenOpts().FP32DenormalMode = llvm::DenormalMode::getDynamic(); + // FP32 denormal mode + switch (clspv::Option::ExecutionModeDenorm( + clspv::Option::FloatingPointType::fp32)) { + case clspv::Option::DenormMode::preserve: + instance.getCodeGenOpts().FP32DenormalMode = llvm::DenormalMode::getIEEE(); + break; + case clspv::Option::DenormMode::flush_to_zero: + instance.getCodeGenOpts().FP32DenormalMode = + llvm::DenormalMode::getPreserveSign(); + break; + case clspv::Option::DenormMode::unspecified: + instance.getCodeGenOpts().FP32DenormalMode = + llvm::DenormalMode::getDynamic(); + break; + case clspv::Option::DenormMode::error: + llvm_unreachable("Invalid DenormMode"); + } + + // FPDenormalMode (for 16 and 64) + clspv::Option::DenormMode mode = clspv::Option::DenormMode::error; + auto modefp16 = clspv::Option::ExecutionModeDenorm( + clspv::Option::FloatingPointType::fp16); + auto modefp64 = clspv::Option::ExecutionModeDenorm( + clspv::Option::FloatingPointType::fp64); + if (modefp16 == modefp64) { + mode = modefp16; + } else if (modefp16 == clspv::Option::DenormMode::unspecified) { + mode = modefp64; + } else if (modefp64 == clspv::Option::DenormMode::unspecified) { + mode = modefp16; + } + switch (mode) { + case clspv::Option::DenormMode::preserve: + instance.getCodeGenOpts().FPDenormalMode = llvm::DenormalMode::getIEEE(); + break; + case clspv::Option::DenormMode::flush_to_zero: + instance.getCodeGenOpts().FPDenormalMode = + llvm::DenormalMode::getPreserveSign(); + break; + case clspv::Option::DenormMode::unspecified: + instance.getCodeGenOpts().FPDenormalMode = llvm::DenormalMode::getDynamic(); + break; + case clspv::Option::DenormMode::error: + llvm_unreachable("Invalid DenormMode"); + } + instance.getCodeGenOpts().StackRealignment = true; instance.getCodeGenOpts().SimplifyLibCalls = false; instance.getCodeGenOpts().EmitOpenCLArgMetadata = false; @@ -341,7 +387,6 @@ int SetCompilerInstanceOptions( instance.getLangOpts().SinglePrecisionConstants = cl_single_precision_constants; - // cl_denorms_are_zero ignored for now! // cl_fp32_correctly_rounded_divide_sqrt ignored for now! instance.getCodeGenOpts().LessPreciseFPMAD = clspv::Option::ClMadEnable() || clspv::Option::UnsafeMath(); @@ -965,13 +1010,13 @@ int ParseOptions(const int argc, const char *const argv[]) { } if (!clspv::Option::FP16() && - ExecutionModeRoundingModeRTE(clspv::Option::RoundingModeRTE::fp16)) { + ExecutionModeRoundingModeRTE(clspv::Option::FloatingPointType::fp16)) { llvm::errs() << "error: Cannot set RoundingModeRTE for fp16 if fp16 is not " "supported\n"; return -1; } if (!clspv::Option::FP64() && - ExecutionModeRoundingModeRTE(clspv::Option::RoundingModeRTE::fp64)) { + ExecutionModeRoundingModeRTE(clspv::Option::FloatingPointType::fp64)) { llvm::errs() << "error: Cannot set RoundingModeRTE for fp64 if fp64 is not " "supported\n"; return -1; @@ -1001,6 +1046,33 @@ int ParseOptions(const int argc, const char *const argv[]) { clspv::Option::AddUseNativeBuiltins(clspv::Builtins::BuiltinType::kFma); } + if (clspv::Option::FP16() && + ExecutionModeDenorm(clspv::Option::FloatingPointType::fp16) == + clspv::Option::DenormMode::error) { + llvm::errs() + << "error: denorm preserve & flush to zero mode are exclusive (fp16)"; + return -1; + } + if (ExecutionModeDenorm(clspv::Option::FloatingPointType::fp32) == + clspv::Option::DenormMode::error) { + llvm::errs() + << "error: denorm preserve & flush to zero mode are exclusive (fp32)"; + return -1; + } + if (clspv::Option::FP64() && + ExecutionModeDenorm(clspv::Option::FloatingPointType::fp64) == + clspv::Option::DenormMode::error) { + llvm::errs() + << "error: denorm preserve & flush to zero mode are exclusive (fp64)"; + return -1; + } + if (clspv::Option::FP16() && clspv::Option::FP64() && + (ExecutionModeDenorm(clspv::Option::FloatingPointType::fp16) != + ExecutionModeDenorm(clspv::Option::FloatingPointType::fp64))) { + llvm::errs() << "error: fp16 & fp64 needs to share the same DenormMode"; + return -1; + } + return 0; } diff --git a/lib/Option.cpp b/lib/Option.cpp index e95cf0bbd..7d86a971d 100644 --- a/lib/Option.cpp +++ b/lib/Option.cpp @@ -296,19 +296,46 @@ static llvm::cl::opt cluster_non_pointer_kernel_args( "other arguments. Use this to reduce storage buffer " "descriptors.")); -static llvm::cl::list rounding_mode_rte( +static llvm::cl::list rounding_mode_rte( "rounding-mode-rte", llvm::cl::desc( "Set execution mode RoundingModeRTE for a floating point type"), llvm::cl::CommaSeparated, llvm::cl::ZeroOrMore, - llvm::cl::values(clEnumValN(clspv::Option::RoundingModeRTE::fp16, "16", + llvm::cl::values(clEnumValN(clspv::Option::FloatingPointType::fp16, "16", "Set execution mode RoundingModeRTE for fp16")), - llvm::cl::values(clEnumValN(clspv::Option::RoundingModeRTE::fp32, "32", + llvm::cl::values(clEnumValN(clspv::Option::FloatingPointType::fp32, "32", "Set execution mode RoundingModeRTE for fp32")), llvm::cl::values( - clEnumValN(clspv::Option::RoundingModeRTE::fp64, "64", + clEnumValN(clspv::Option::FloatingPointType::fp64, "64", "Set execution mode RoundingModeRTE for fp64"))); +static llvm::cl::list denorm_preserve( + "denorm-preserve", + llvm::cl::desc( + "Set execution mode DenormPreserve for a floating point type"), + llvm::cl::CommaSeparated, llvm::cl::ZeroOrMore, + llvm::cl::values(clEnumValN(clspv::Option::FloatingPointType::fp16, "16", + "Set execution mode DenormPreserve for fp16")), + llvm::cl::values(clEnumValN(clspv::Option::FloatingPointType::fp32, "32", + "Set execution mode DenormPreserve for fp32")), + llvm::cl::values(clEnumValN(clspv::Option::FloatingPointType::fp64, "64", + "Set execution mode DenormPreserve for fp64"))); + +static llvm::cl::list denorm_flush_to_zero( + "denorm-flush-to-zero", + llvm::cl::desc( + "Set execution mode DenormFlushToZero for a floating point type"), + llvm::cl::CommaSeparated, llvm::cl::ZeroOrMore, + llvm::cl::values( + clEnumValN(clspv::Option::FloatingPointType::fp16, "16", + "Set execution mode DenormFlushToZero for fp16")), + llvm::cl::values( + clEnumValN(clspv::Option::FloatingPointType::fp32, "32", + "Set execution mode DenormFlushToZero for fp32")), + llvm::cl::values( + clEnumValN(clspv::Option::FloatingPointType::fp64, "64", + "Set execution mode DenormFlushToZero for fp64"))); + static llvm::cl::list no_16bit_storage( "no-16bit-storage", llvm::cl::desc("Disable fine-grained 16-bit storage capabilities."), @@ -453,15 +480,15 @@ static llvm::cl::opt untyped_pointers("untyped-pointers", llvm::cl::init(false), llvm::cl::desc("Enable SPV_KHR_untyped_pointers")); -static llvm::cl::list spv_khr_fma( +static llvm::cl::list spv_khr_fma( "spv-khr-fma", llvm::cl::desc("Enable SPV_KHR_fma for a floating point type"), llvm::cl::CommaSeparated, llvm::cl::ZeroOrMore, - llvm::cl::values(clEnumValN(clspv::Option::SpvKhrFma::fp16, "16", + llvm::cl::values(clEnumValN(clspv::Option::FloatingPointType::fp16, "16", "Enable SPV_KHR_fma for fp16")), - llvm::cl::values(clEnumValN(clspv::Option::SpvKhrFma::fp32, "32", + llvm::cl::values(clEnumValN(clspv::Option::FloatingPointType::fp32, "32", "Enable SPV_KHR_fma for fp32")), - llvm::cl::values(clEnumValN(clspv::Option::SpvKhrFma::fp64, "64", + llvm::cl::values(clEnumValN(clspv::Option::FloatingPointType::fp64, "64", "Enable SPV_KHR_fma for fp64"))); } // namespace @@ -518,7 +545,7 @@ bool NonUniformNDRangeSupported() { } bool ClusterPodKernelArgs() { return cluster_non_pointer_kernel_args; } -bool ExecutionModeRoundingModeRTE(RoundingModeRTE fp) { +bool ExecutionModeRoundingModeRTE(FloatingPointType fp) { for (auto type : rounding_mode_rte) { if (type == fp) { return true; @@ -527,6 +554,32 @@ bool ExecutionModeRoundingModeRTE(RoundingModeRTE fp) { return false; } +DenormMode ExecutionModeDenorm(FloatingPointType fpty) { + static std::unordered_map modes; + auto mode = modes.find(fpty); + if (mode != modes.end()) { + return modes[fpty]; + } + + for (auto type : denorm_preserve) { + modes[type] = DenormMode::preserve; + } + for (auto type : denorm_flush_to_zero) { + if (modes.find(type) != modes.end()) { + modes[type] = DenormMode::error; + } else { + modes[type] = DenormMode::flush_to_zero; + } + } + for (auto type : {FloatingPointType::fp16, FloatingPointType::fp32, + FloatingPointType::fp64}) { + if (modes.find(type) == modes.end()) { + modes[type] = DenormMode::unspecified; + } + } + return modes[fpty]; +} + bool Supports16BitStorageClass(StorageClass sc) { // -no-16bit-storage removes storage capabilities. for (auto storage_class : no_16bit_storage) { @@ -625,16 +678,16 @@ bool UntypedPointerAddressSpace(unsigned aspace) { } bool SupportsFmaKHR(uint32_t scalarSizeInBits) { - SpvKhrFma fma; + FloatingPointType fma; switch (scalarSizeInBits) { case 16: - fma = clspv::Option::SpvKhrFma::fp16; + fma = clspv::Option::FloatingPointType::fp16; break; case 32: - fma = clspv::Option::SpvKhrFma::fp32; + fma = clspv::Option::FloatingPointType::fp32; break; case 64: - fma = clspv::Option::SpvKhrFma::fp64; + fma = clspv::Option::FloatingPointType::fp64; break; default: return false; diff --git a/lib/SPIRVProducerPass.cpp b/lib/SPIRVProducerPass.cpp index adfad3496..b0bc34b97 100644 --- a/lib/SPIRVProducerPass.cpp +++ b/lib/SPIRVProducerPass.cpp @@ -387,9 +387,9 @@ struct SPIRVProducerPassImpl { bool hasConvertToF() { return HasConvertToF; } void setConvertToF() { if (!HasConvertToF && - (ExecutionModeRoundingModeRTE(RoundingModeRTE::fp16) || - ExecutionModeRoundingModeRTE(RoundingModeRTE::fp32) || - ExecutionModeRoundingModeRTE(RoundingModeRTE::fp64))) { + (ExecutionModeRoundingModeRTE(FloatingPointType::fp16) || + ExecutionModeRoundingModeRTE(FloatingPointType::fp32) || + ExecutionModeRoundingModeRTE(FloatingPointType::fp64))) { addCapability(spv::CapabilityRoundingModeRTE); HasConvertToF = true; } @@ -3377,6 +3377,21 @@ void SPIRVProducerPassImpl::GenerateModuleInfo() { auto &EntryPointInterfaces = getEntryPointInterfacesList(); std::vector &BuiltinDimVec = getBuiltinDimVec(); + // Add denorm capabilities if needed + if ((ExecutionModeDenorm(FloatingPointType::fp16) == DenormMode::preserve) || + (ExecutionModeDenorm(FloatingPointType::fp32) == DenormMode::preserve) || + (ExecutionModeDenorm(FloatingPointType::fp64) == DenormMode::preserve)) { + addCapability(spv::CapabilityDenormPreserve); + } + if ((ExecutionModeDenorm(FloatingPointType::fp16) == + DenormMode::flush_to_zero) || + (ExecutionModeDenorm(FloatingPointType::fp32) == + DenormMode::flush_to_zero) || + (ExecutionModeDenorm(FloatingPointType::fp64) == + DenormMode::flush_to_zero)) { + addCapability(spv::CapabilityDenormFlushToZero); + } + SPIRVOperandVec Ops; for (auto Capability : CapabilitySet) { @@ -3408,7 +3423,16 @@ void SPIRVProducerPassImpl::GenerateModuleInfo() { } } - if (SpvVersion() < SPIRVVersion::SPIRV_1_4 && hasConvertToF()) { + bool hasDenorms = + (clspv::Option::ExecutionModeDenorm(FloatingPointType::fp16) != + DenormMode::unspecified) || + (clspv::Option::ExecutionModeDenorm(FloatingPointType::fp32) != + DenormMode::unspecified) || + (clspv::Option::ExecutionModeDenorm(FloatingPointType::fp64) != + DenormMode::unspecified); + + if (SpvVersion() < SPIRVVersion::SPIRV_1_4 && + (hasConvertToF() || hasDenorms)) { addSPIRVInst(spv::OpExtension, "SPV_KHR_float_controls"); } @@ -3567,18 +3591,18 @@ void SPIRVProducerPassImpl::GenerateModuleInfo() { if (hasConvertToF()) { for (auto EntryPoint : EntryPoints) { if (clspv::Option::FP16() && - ExecutionModeRoundingModeRTE(RoundingModeRTE::fp16)) { + ExecutionModeRoundingModeRTE(FloatingPointType::fp16)) { Ops.clear(); Ops << EntryPoint.second << spv::ExecutionModeRoundingModeRTE << 16; addSPIRVInst(spv::OpExecutionMode, Ops); } - if (ExecutionModeRoundingModeRTE(RoundingModeRTE::fp32)) { + if (ExecutionModeRoundingModeRTE(FloatingPointType::fp32)) { Ops.clear(); Ops << EntryPoint.second << spv::ExecutionModeRoundingModeRTE << 32; addSPIRVInst(spv::OpExecutionMode, Ops); } if (clspv::Option::FP64() && - ExecutionModeRoundingModeRTE(RoundingModeRTE::fp64)) { + ExecutionModeRoundingModeRTE(FloatingPointType::fp64)) { Ops.clear(); Ops << EntryPoint.second << spv::ExecutionModeRoundingModeRTE << 64; addSPIRVInst(spv::OpExecutionMode, Ops); @@ -3586,6 +3610,40 @@ void SPIRVProducerPassImpl::GenerateModuleInfo() { } } + // Emit denorm execution modes + for (auto EntryPoint : EntryPoints) { + auto mode = ExecutionModeDenorm(FloatingPointType::fp16); + if (FP16() && mode == DenormMode::preserve) { + Ops.clear(); + Ops << EntryPoint.second << spv::ExecutionModeDenormPreserve << 16; + addSPIRVInst(spv::OpExecutionMode, Ops); + } else if (FP16() && mode == DenormMode::flush_to_zero) { + Ops.clear(); + Ops << EntryPoint.second << spv::ExecutionModeDenormFlushToZero << 16; + addSPIRVInst(spv::OpExecutionMode, Ops); + } + mode = ExecutionModeDenorm(FloatingPointType::fp32); + if (mode == DenormMode::preserve) { + Ops.clear(); + Ops << EntryPoint.second << spv::ExecutionModeDenormPreserve << 32; + addSPIRVInst(spv::OpExecutionMode, Ops); + } else if (mode == DenormMode::flush_to_zero) { + Ops.clear(); + Ops << EntryPoint.second << spv::ExecutionModeDenormFlushToZero << 32; + addSPIRVInst(spv::OpExecutionMode, Ops); + } + mode = ExecutionModeDenorm(FloatingPointType::fp64); + if (FP64() && mode == DenormMode::preserve) { + Ops.clear(); + Ops << EntryPoint.second << spv::ExecutionModeDenormPreserve << 64; + addSPIRVInst(spv::OpExecutionMode, Ops); + } else if (FP64() && mode == DenormMode::flush_to_zero) { + Ops.clear(); + Ops << EntryPoint.second << spv::ExecutionModeDenormFlushToZero << 64; + addSPIRVInst(spv::OpExecutionMode, Ops); + } + } + // // Generate OpSource. // diff --git a/test/denorm_modes.cl b/test/denorm_modes.cl new file mode 100644 index 000000000..c43d118e9 --- /dev/null +++ b/test/denorm_modes.cl @@ -0,0 +1,49 @@ +// RUN: clspv %s -o %t.preserve.spv -denorm-preserve=16,32,64 +// RUN: spirv-dis %t.preserve.spv -o %t.preserve.spvasm +// RUN: FileCheck %s --check-prefix=PRESERVE < %t.preserve.spvasm +// RUN: spirv-val --target-env spv1.4 %t.preserve.spv + +// RUN: clspv %s -o %t.flush.spv -denorm-flush-to-zero=16,32,64 +// RUN: spirv-dis %t.flush.spv -o %t.flush.spvasm +// RUN: FileCheck %s --check-prefix=FLUSH < %t.flush.spvasm +// RUN: spirv-val --target-env spv1.4 %t.flush.spv + +// RUN: clspv %s -o %t.none.spv +// RUN: spirv-dis %t.none.spv -o %t.none.spvasm +// RUN: FileCheck %s --check-prefix=NONE < %t.none.spvasm +// RUN: spirv-val %t.none.spv + +// RUN: clspv %s -o %t.preserve14.spv -denorm-preserve=16,32,64 -spv-version 1.4 +// RUN: spirv-dis %t.preserve14.spv -o %t.preserve14.spvasm +// RUN: FileCheck %s --check-prefix=PRESERVE14 < %t.preserve14.spvasm +// RUN: spirv-val --target-env spv1.4 %t.preserve14.spv + +// PRESERVE: OpCapability DenormPreserve +// PRESERVE: OpExtension "SPV_KHR_float_controls" +// PRESERVE: OpExecutionMode {{.*}} DenormPreserve 16 +// PRESERVE: OpExecutionMode {{.*}} DenormPreserve 32 +// PRESERVE: OpExecutionMode {{.*}} DenormPreserve 64 + +// FLUSH: OpCapability DenormFlushToZero +// FLUSH: OpExtension "SPV_KHR_float_controls" +// FLUSH: OpExecutionMode {{.*}} DenormFlushToZero 16 +// FLUSH: OpExecutionMode {{.*}} DenormFlushToZero 32 +// FLUSH: OpExecutionMode {{.*}} DenormFlushToZero 64 + +// NONE-NOT: OpCapability DenormPreserve +// NONE-NOT: OpCapability DenormFlushToZero +// NONE-NOT: OpExtension "SPV_KHR_float_controls" +// NONE-NOT: OpExecutionMode {{.*}} DenormPreserve +// NONE-NOT: OpExecutionMode {{.*}} DenormFlushToZero + +// PRESERVE14: OpCapability DenormPreserve +// PRESERVE14-NOT: OpExtension "SPV_KHR_float_controls" +// PRESERVE14: OpExecutionMode {{.*}} DenormPreserve 16 +// PRESERVE14: OpExecutionMode {{.*}} DenormPreserve 32 +// PRESERVE14: OpExecutionMode {{.*}} DenormPreserve 64 + +void kernel foo(global int *input, global float *output) +{ + uint gid = get_global_id(0); + output[gid] = input[gid]; +}