diff --git a/examples/rhi/floattexture/floattexture.cpp b/examples/rhi/floattexture/floattexture.cpp
index 1ef696c80b6846b60486f6674de173977b227f1f..021767ff212dd9fc2c36ec3b27751040a1ef09dd 100644
--- a/examples/rhi/floattexture/floattexture.cpp
+++ b/examples/rhi/floattexture/floattexture.cpp
@@ -286,7 +286,7 @@ void Window::customInit()
     d.initialUpdates->uploadStaticBuffer(d.vbuf, vertexData);
     d.initialUpdates->uploadStaticBuffer(d.ibuf, indexData);
 
-    qint32 flip = m_r->isYUpInFramebuffer() ? 0 : 1;
+    qint32 flip = 1;
     d.initialUpdates->updateDynamicBuffer(d.ubuf, 64, 4, &flip);
 
     QRhiTextureMipLevel mipDesc(floatData);
diff --git a/src/rhi/qrhigles2.cpp b/src/rhi/qrhigles2.cpp
index fb48e3edb827a260ba4f5589a89fcdf4c9974ec0..839d5a8eedcf92162fac51a79a7c364b5f4ca682 100644
--- a/src/rhi/qrhigles2.cpp
+++ b/src/rhi/qrhigles2.cpp
@@ -941,6 +941,7 @@ void QRhiGles2::enqueueResourceUpdates(QRhiCommandBuffer *cb, QRhiResourceUpdate
         if (u.type == QRhiResourceUpdateBatchPrivate::TextureOp::TexUpload) {
             QGles2Texture *texD = QRHI_RES(QGles2Texture, u.upload.tex);
             const bool isCompressed = isCompressedFormat(texD->m_format);
+            const bool isFloat = isFloatFormat(texD->m_format);
             const bool isCubeMap = texD->m_flags.testFlag(QRhiTexture::CubeMap);
             const GLenum faceTargetBase = isCubeMap ? GL_TEXTURE_CUBE_MAP_POSITIVE_X : texD->target;
             const QVector<QRhiTextureLayer> layers = u.upload.desc.layers();
@@ -983,7 +984,24 @@ void QRhiGles2::enqueueResourceUpdates(QRhiCommandBuffer *cb, QRhiResourceUpdate
                             cmd.args.compressedImage.data = cbD->retainData(rawData);
                             cbD->commands.append(cmd);
                         }
-                    } else {
+                    } else if (isFloat && !rawData.isEmpty()) {
+                        const QSize size = mipDesc.sourceSize().isEmpty() ? q->sizeForMipLevel(level, texD->m_pixelSize)
+                                                                          : mipDesc.sourceSize();
+                        QGles2CommandBuffer::Command cmd;
+                        cmd.cmd = QGles2CommandBuffer::Command::SubImage;
+                        cmd.args.subImage.target = texD->target;
+                        cmd.args.subImage.texture = texD->texture;
+                        cmd.args.subImage.faceTarget = faceTargetBase + layer;
+                        cmd.args.subImage.level = level;
+                        cmd.args.subImage.dx = dp.x();
+                        cmd.args.subImage.dy = dp.y();
+                        cmd.args.subImage.w = size.width();
+                        cmd.args.subImage.h = size.height();
+                        cmd.args.subImage.glformat = texD->glformat;
+                        cmd.args.subImage.gltype = texD->gltype;
+                        cmd.args.subImage.data = cbD->retainData(rawData);
+                        cbD->commands.append(cmd);
+                    } else if (!mipDesc.image().isNull()) {
                         QImage img = mipDesc.image();
                         QSize size = img.size();
                         QGles2CommandBuffer::Command cmd;
@@ -1006,6 +1024,8 @@ void QRhiGles2::enqueueResourceUpdates(QRhiCommandBuffer *cb, QRhiResourceUpdate
                         cmd.args.subImage.gltype = texD->gltype;
                         cmd.args.subImage.data = cbD->retainImage(img);
                         cbD->commands.append(cmd);
+                    } else {
+                        qWarning("Invalid texture upload for %p layer=%d mip=%d", texD, layer, level);
                     }
                 }
             }
@@ -2063,10 +2083,12 @@ bool QGles2Texture::prepareBuild(QSize *adjustedSize)
         case QRhiTexture::RGBA16F:
             glintformat = GL_RGBA16F;
             glformat = GL_RGBA;
+            gltype = GL_FLOAT;
             break;
         case QRhiTexture::RGBA32F:
             glintformat = GL_RGBA32F;
             glformat = GL_RGBA;
+            gltype = GL_FLOAT;
             break;
         default:
             Q_UNREACHABLE();
diff --git a/src/rhi/qrhivulkan.cpp b/src/rhi/qrhivulkan.cpp
index f0647b2f48298330e23f8d2b75a5600ccd664bfa..bca73f480dc922bfbeefe2dd92ca1c42ca0f7c45 100644
--- a/src/rhi/qrhivulkan.cpp
+++ b/src/rhi/qrhivulkan.cpp
@@ -2396,7 +2396,22 @@ void QRhiVulkan::enqueueResourceUpdates(QRhiCommandBuffer *cb, QRhiResourceUpdat
                         copyInfo.imageExtent.width = size.width();
                         copyInfo.imageExtent.height = size.height();
                         copyInfos.append(copyInfo);
-                    } else if (!rawData.isEmpty()) {
+                    } else if (!rawData.isEmpty() && isFloatFormat(utexD->m_format)) {
+                        copySizeBytes = imageSizeBytes = rawData.size();
+                        src = rawData.constData();
+                        QSize size = q->sizeForMipLevel(level, utexD->m_pixelSize);
+                        const int subresw = size.width();
+                        const int subresh = size.height();
+                        if (!mipDesc.sourceSize().isEmpty())
+                            size = mipDesc.sourceSize();
+                        const int w = size.width();
+                        const int h = size.height();
+                        copyInfo.imageOffset.x = dp.x();
+                        copyInfo.imageOffset.y = dp.y();
+                        copyInfo.imageExtent.width = w;
+                        copyInfo.imageExtent.height = h;
+                        copyInfos.append(copyInfo);
+                    } else if (!rawData.isEmpty() && isCompressedFormat(utexD->m_format)) {
                         copySizeBytes = imageSizeBytes = rawData.size();
                         src = rawData.constData();
                         QSize size = q->sizeForMipLevel(level, utexD->m_pixelSize);
@@ -2416,6 +2431,8 @@ void QRhiVulkan::enqueueResourceUpdates(QRhiCommandBuffer *cb, QRhiResourceUpdat
                         copyInfo.imageExtent.width = dp.x() + w == subresw ? w : aligned(w, blockDim.width());
                         copyInfo.imageExtent.height = dp.y() + h == subresh ? h : aligned(h, blockDim.height());
                         copyInfos.append(copyInfo);
+                    } else {
+                        qWarning("Invalid texture upload for %p layer=%d mip=%d", utexD, layer, level);
                     }
 
                     memcpy(reinterpret_cast<char *>(mp) + curOfs, src, copySizeBytes);