initial version

samples/images/10_mipmapimage/CMakeLists.txt

@@ -0,0 +1,11 @@
+# Copyright (c) 2024 Ben Ashbaugh
+#
+# SPDX-License-Identifier: MIT
+
+add_opencl_sample(
+    TEST
+    NUMBER 10
+    TARGET mipmapimage
+    VERSION 300
+    CATEGORY images
+    SOURCES main.cpp)

samples/images/10_mipmapimage/main.cpp

@@ -0,0 +1,185 @@
+/*
+// Copyright (c) 2019-2024 Ben Ashbaugh
+//
+// SPDX-License-Identifier: MIT
+*/
+
+#include <popl/popl.hpp>
+
+#include <CL/opencl.hpp>
+
+#include "util.hpp"
+
+#include <algorithm>
+
+static const char kernelString[] = R"CLC(
+kernel void CopyBuffer( global uint* dst, global uint* src )
+{
+    uint id = get_global_id(0);
+    dst[id] = src[id];
+}
+)CLC";
+
+int main(
+    int argc,
+    char** argv )
+{
+    int platformIndex = 0;
+    int deviceIndex = 0;
+
+    cl_uint mipLevels = 6;
+
+    {
+        popl::OptionParser op("Supported Options");
+        op.add<popl::Value<int>>("p", "platform", "Platform Index", platformIndex, &platformIndex);
+        op.add<popl::Value<int>>("d", "device", "Device Index", deviceIndex, &deviceIndex);
+        op.add<popl::Value<cl_uint>>("m", "miplevels", "Number of Mipmap Levels", mipLevels, &mipLevels);
+
+        bool printUsage = false;
+        try {
+            op.parse(argc, argv);
+        } catch (std::exception& e) {
+            fprintf(stderr, "Error: %s\n\n", e.what());
+            printUsage = true;
+        }
+
+        if (printUsage || !op.unknown_options().empty() || !op.non_option_args().empty()) {
+            fprintf(stderr,
+                "Usage: copybufferkernel [options]\n"
+                "%s", op.help().c_str());
+            return -1;
+        }
+    }
+
+    std::vector<cl::Platform> platforms;
+    cl::Platform::get(&platforms);
+
+    printf("Running on platform: %s\n",
+        platforms[platformIndex].getInfo<CL_PLATFORM_NAME>().c_str() );
+
+    std::vector<cl::Device> devices;
+    platforms[platformIndex].getDevices(CL_DEVICE_TYPE_ALL, &devices);
+
+    printf("Running on device: %s\n",
+        devices[deviceIndex].getInfo<CL_DEVICE_NAME>().c_str() );
+
+    bool has_cl_khr_mipmap_image =
+        checkDeviceForExtension(devices[deviceIndex], CL_KHR_MIPMAP_IMAGE_EXTENSION_NAME);
+    if (has_cl_khr_mipmap_image) {
+        printf("Device supports " CL_KHR_MIPMAP_IMAGE_EXTENSION_NAME ".\n");
+    } else {
+        printf("Device does not support " CL_KHR_MIPMAP_IMAGE_EXTENSION_NAME ", exiting.\n");
+        return -1;
+    }
+
+    cl::Context context{devices[deviceIndex]};
+    cl::CommandQueue commandQueue{context, devices[deviceIndex]};
+
+    cl::Program program{ context, kernelString };
+    program.build();
+    cl::Kernel kernel = cl::Kernel{ program, "CopyBuffer" };
+
+    if (mipLevels > 20) {
+        printf("Maximum number of supported mip levels is 20.\n");
+        mipLevels = 20;
+    }
+
+    size_t imageWidth = std::max(32, 1 << mipLevels);
+    size_t imageHeight = std::max(32, 1 << mipLevels);
+
+    cl_image_format imgFormat{};
+    imgFormat.image_channel_order = CL_R;
+    imgFormat.image_channel_data_type = CL_FLOAT;
+
+    cl_image_desc imgDesc{};
+    imgDesc.image_type = CL_MEM_OBJECT_IMAGE2D;
+    imgDesc.image_width = imageWidth;
+    imgDesc.image_height = imageHeight;
+    imgDesc.num_mip_levels = mipLevels;
+
+    printf("Creating a %zu x %zu image with %u mip levels...\n",
+        imgDesc.image_width, imgDesc.image_height, imgDesc.num_mip_levels);
+
+    cl::Image2D img = cl::Image2D{
+        clCreateImage(
+            context(),
+            CL_MEM_READ_ONLY,
+            &imgFormat,
+            &imgDesc,
+            nullptr,
+            nullptr)};
+
+    cl::Buffer buf = cl::Buffer{
+        context,
+        CL_MEM_ALLOC_HOST_PTR,
+        64 * sizeof(cl_float)};
+
+    // initialization
+    {
+        printf("Initializing mipmap image...\n");
+        for (cl_uint m = 0; m < mipLevels; m++) {
+            const cl_float v = (m + 1.0f) / mipLevels;
+            const cl_float4 mipData{v, v, v, v};
+            commandQueue.enqueueFillImage(
+                img,
+                mipData,
+                {0, 0, m},
+                {imageWidth >> m, imageHeight >> m, 1});
+        }
+        commandQueue.finish();
+    }
+
+#if 0
+    // execution
+    kernel.setArg(0, deviceMemDst);
+    kernel.setArg(1, deviceMemSrc);
+    commandQueue.enqueueNDRangeKernel(
+        kernel,
+        cl::NullRange,
+        cl::NDRange{gwx} );
+
+    // verification
+    {
+        const cl_uint*  pDst = (const cl_uint*)commandQueue.enqueueMapBuffer(
+            deviceMemDst,
+            CL_TRUE,
+            CL_MAP_READ,
+            0,
+            gwx * sizeof(cl_uint) );
+
+        unsigned int    mismatches = 0;
+
+        for( size_t i = 0; i < gwx; i++ )
+        {
+            if( pDst[i] != i )
+            {
+                if( mismatches < 16 )
+                {
+                    fprintf(stderr, "MisMatch!  dst[%d] == %08X, want %08X\n",
+                        (unsigned int)i,
+                        pDst[i],
+                        (unsigned int)i );
+                }
+                mismatches++;
+            }
+        }
+
+        if( mismatches )
+        {
+            fprintf(stderr, "Error: Found %d mismatches / %d values!!!\n",
+                mismatches,
+                (unsigned int)gwx );
+        }
+        else
+        {
+            printf("Success.\n");
+        }
+
+        commandQueue.enqueueUnmapMemObject(
+            deviceMemDst,
+            (void*)pDst );
+    }
+#endif
+
+    return 0;
+}

samples/images/CMakeLists.txt

@@ -2,7 +2,5 @@
 #
 # SPDX-License-Identifier: MIT
 
-# Require OpenCL 2.0, for read-write images:
-set(SAMPLES_CL_VERSION 200)
-
 add_subdirectory( 00_enumimageformats )
+add_subdirectory( 10_mipmapimage )