diff options
| author | jsmall-nvidia <jsmall@nvidia.com> | 2021-04-23 11:32:07 -0400 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2021-04-23 11:32:07 -0400 |
| commit | 79e722338cd59aab74b4c57600c5ac6bce3bcd25 (patch) | |
| tree | af77066235c4038bd15c6297ef4f48d3e562171d /source/slang/slang-emit-cuda.cpp | |
| parent | a47e7751c2738543e872452debc7494369c9fb35 (diff) | |
Preliminary CUDA Half support (#1808)
* #include an absolute path didn't work - because paths were taken to always be relative.
* WIP CUDA half support.
* Working support for half on CUDA - requires cuda_fp16.h and associated files can be found.
* Fix for win32 for unused funcs.
* Fix for Clang.
* Hack to disable unused local function warning.
Diffstat (limited to 'source/slang/slang-emit-cuda.cpp')
| -rw-r--r-- | source/slang/slang-emit-cuda.cpp | 24 |
1 files changed, 12 insertions, 12 deletions
diff --git a/source/slang/slang-emit-cuda.cpp b/source/slang/slang-emit-cuda.cpp index 2f5a9917d..a259ea933 100644 --- a/source/slang/slang-emit-cuda.cpp +++ b/source/slang/slang-emit-cuda.cpp @@ -27,7 +27,7 @@ static bool _isSingleNameBasicType(IROp op) } } -/* static */ UnownedStringSlice CUDASourceEmitter::getBuiltinTypeName(IROp op) +UnownedStringSlice CUDASourceEmitter::getBuiltinTypeName(IROp op) { switch (op) { @@ -44,8 +44,11 @@ static bool _isSingleNameBasicType(IROp op) case kIROp_UIntType: return UnownedStringSlice("uint"); case kIROp_UInt64Type: return UnownedStringSlice("ulonglong"); - // Not clear just yet how we should handle half... we want all processing as float probly, but when reading/writing to memory converting - case kIROp_HalfType: return UnownedStringSlice("half"); + case kIROp_HalfType: + { + m_extensionTracker->requireBaseType(BaseType::Half); + return UnownedStringSlice("__half"); + } case kIROp_FloatType: return UnownedStringSlice("float"); case kIROp_DoubleType: return UnownedStringSlice("double"); @@ -54,7 +57,7 @@ static bool _isSingleNameBasicType(IROp op) } -/* static */ UnownedStringSlice CUDASourceEmitter::getVectorPrefix(IROp op) +UnownedStringSlice CUDASourceEmitter::getVectorPrefix(IROp op) { switch (op) { @@ -70,8 +73,11 @@ static bool _isSingleNameBasicType(IROp op) case kIROp_UIntType: return UnownedStringSlice("uint"); case kIROp_UInt64Type: return UnownedStringSlice("ulonglong"); - // Not clear just yet how we should handle half... we want all processing as float probly, but when reading/writing to memory converting - case kIROp_HalfType: return UnownedStringSlice("half"); + case kIROp_HalfType: + { + m_extensionTracker->requireBaseType(BaseType::Half); + return UnownedStringSlice("__half"); + } case kIROp_FloatType: return UnownedStringSlice("float"); case kIROp_DoubleType: return UnownedStringSlice("double"); @@ -160,12 +166,6 @@ SlangResult CUDASourceEmitter::calcTypeName(IRType* type, CodeGenTarget target, switch (type->getOp()) { - case kIROp_HalfType: - { - // Special case half - out << getBuiltinTypeName(kIROp_FloatType); - return SLANG_OK; - } case kIROp_VectorType: { auto vecType = static_cast<IRVectorType*>(type); |