diff options
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); |