Fixed issues in ImageInverter CL code

Author: smistad

source/FAST/Algorithms/ImageInverter/ImageInverter.cpp

@@ -31,23 +31,15 @@ void ImageInverter::execute() {
 
     if(input->getDimensions() == 3) {
         std::string buildOptions = "-DDATA_TYPE=" + getCTypeAsString(output->getDataType());
-        cl::Program program = getOpenCLProgram(device, "3D", buildOptions);
-        cl::Kernel kernel(program, "invert3D");
+        auto kernel = getKernel("invert3D", "3D", buildOptions);
 
-        auto access = input->getOpenCLImageAccess(ACCESS_READ, device);
-        auto access2 = output->getOpenCLBufferAccess(ACCESS_READ_WRITE, device);
-        kernel.setArg(0, *access->get3DImage());
-        kernel.setArg(1, *access2->get());
-        kernel.setArg(2, min);
-        kernel.setArg(3, max);
-        kernel.setArg(4, output->getNrOfChannels());
+        kernel.setArg("input", input);
+        kernel.setArg("output", output);
+        kernel.setArg("minIntensity", min);
+        kernel.setArg("maxIntensity", max);
+        kernel.setArg("outputChannels", output->getNrOfChannels());
 
-        queue.enqueueNDRangeKernel(
-                kernel,
-                cl::NullRange,
-                cl::NDRange(size.x(), size.y(), size.z()),
-                cl::NullRange
-        );
+        getQueue().add(kernel, size);
     } else {
         auto kernel = getKernel("invert2D", "2D");
 

source/FAST/Algorithms/ImageInverter/ImageInverter2D.cl

@@ -17,7 +17,7 @@ __kernel void invert2D(
     } else {
         value = convert_float4(read_imagei(input, sampler, pos));
     }
-    value = (maxIntensity - minIntensity) - value;
+    value = (maxIntensity + minIntensity) - value;
     if(dataType == CLK_FLOAT) {
         write_imagef(output, pos, value);
     } else if(dataType == CLK_UNSIGNED_INT8 || dataType == CLK_UNSIGNED_INT16 || dataType == CLK_UNSIGNED_INT32) {

source/FAST/Algorithms/ImageInverter/ImageInverter3D.cl

@@ -2,9 +2,9 @@ __constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE |
 
 __kernel void invert3D(
         __read_only image3d_t input,
-        __global DATA_TYPE* output,
-        __private float min,
-        __private float max,
+        __write_only __global DATA_TYPE* output,
+        __private float minIntensity,
+        __private float maxIntensity,
         __private uint outputChannels
         ) {
     const int4 pos = {get_global_id(0), get_global_id(1), get_global_id(2), 0};
@@ -18,7 +18,7 @@ __kernel void invert3D(
     } else {
         value = convert_float4(read_imagei(input, sampler, pos));
     }
-    value = (max - min) - value;
+    value = (maxIntensity + minIntensity) - value;
 
     output[(pos.x + pos.y*get_image_width(input) + pos.z*get_image_width(input)*get_image_height(input))*outputChannels] = value.x;
     if(outputChannels > 1)

source/FAST/Algorithms/ImageInverter/ImageInverterTests.cpp

@@ -8,5 +8,12 @@ TEST_CASE("ImageInverter 2D", "[fast][ImageInverter]") {
     auto importer = ImageFileImporter::create(Config::getTestDataPath() + "US/Heart/ApicalFourChamber/US-2D_0.mhd");
 
     auto inverter = ImageInverter::create()->connect(importer);
-    inverter->run();
-}
+    CHECK_NOTHROW(inverter->run());
+}
+
+TEST_CASE("ImageInverter 3D", "[fast][ImageInverter]") {
+    auto importer = ImageFileImporter::create(Config::getTestDataPath() + "CT/CT-Thorax.mhd");
+
+    auto inverter = ImageInverter::create()->connect(importer);
+    CHECK_NOTHROW(inverter->run());
+}

source/FAST/OpenCLProgram.cpp

@@ -161,7 +161,7 @@ Kernel::Kernel(cl::Kernel clKernel, OpenCLDevice::pointer device) {
     m_device = device;
 
     std::map<cl_kernel_arg_access_qualifier, KernelArgumentAccess> accessMap = {
-            {CL_KERNEL_ARG_ACCESS_NONE, KernelArgumentAccess::NONE},
+            {CL_KERNEL_ARG_ACCESS_NONE, KernelArgumentAccess::UNSPECIFIED},
             {CL_KERNEL_ARG_ACCESS_READ_ONLY, KernelArgumentAccess::READ_ONLY},
             {CL_KERNEL_ARG_ACCESS_WRITE_ONLY, KernelArgumentAccess::WRITE_ONLY},
             {CL_KERNEL_ARG_ACCESS_READ_WRITE, KernelArgumentAccess::READ_WRITE}
@@ -325,10 +325,10 @@ void Kernel::setArg(const std::string& name, OpenCLBuffer buffer) {
 template <>
 void Kernel::setArg(int index, Image::pointer image) {
     checkIndex(index);
-    accessType access = ACCESS_READ;
+    accessType access = ACCESS_READ_WRITE;
     auto kernelAccess = m_argInfoByIndex.at(index).access;
-    if(kernelAccess == KernelArgumentAccess::WRITE_ONLY || kernelAccess == KernelArgumentAccess::READ_WRITE) {
-        access = ACCESS_READ_WRITE;
+    if(kernelAccess == KernelArgumentAccess::READ_ONLY) {
+        access = ACCESS_READ;
     }
     setImageArg(index, image, access);
 }
@@ -342,6 +342,7 @@ template <>
 void Kernel::setArg(int index, std::unique_ptr<OpenCLBufferAccess> access) {
     checkIndex(index);
     m_kernel.setArg(index, *access->get());
+    m_argGotValue.insert(index);
 }
 template <>
 void Kernel::setArg(const std::string& name, std::unique_ptr<OpenCLBufferAccess> access) {

source/FAST/OpenCLProgram.hpp

@@ -46,8 +46,12 @@ class FAST_EXPORT  OpenCLProgram : public Object {
         std::unordered_map<std::shared_ptr<OpenCLDevice>, std::map<std::string, cl::Program> > mOpenCLPrograms;
 };
 
+/**
+ * @brief Access qualifier of a kernel argument
+ * @ingroup opencl
+ */
 enum class KernelArgumentAccess {
-    NONE,
+    UNSPECIFIED,
     READ_ONLY,
     WRITE_ONLY,
     READ_WRITE,
@@ -221,34 +225,86 @@ FAST_EXPORT void Kernel::setArg(const std::string& name, std::unique_ptr<OpenCLB
 class FAST_EXPORT Queue {
     public:
         Queue(cl::CommandQueue clQueue);
+        /**
+         * @brief Enqueue a Kernel to this command queue
+         * @param kernel
+         * @param globalSize total number of work-items for each dimension
+         * @param offset work-item offset, if none is given, it is zero for all dimensions
+         * @param groupSize How many work-items should be each work-group for each dimension.
+         *      If not provided, platform will decide automatically
+         */
         void add(const Kernel& kernel, std::vector<int> globalSize, std::vector<int> offset = std::vector<int>(), std::vector<int> groupSize = std::vector<int>());
+        /**
+         * @brief Block until entire command queue is finished
+         */
         void finish();
+        /**
+         * @brief Copy data from buffer on device to pointer on host
+         * @param buffer OpenCL buffer to read from
+         * @param block Whether this call should block and wait until the data operation is finished.
+         * @param offset Byte offset in OpenCL buffer
+         * @param size nr of bytes to read (e.g. elements*sizeof(datatype))
+         * @param pointerToData Pointer to host memory to read data into. Must be allocated with big enough size
+         */
         void addReadBuffer(OpenCLBuffer buffer, bool block, std::size_t offset, std::size_t size, void* pointerToData);
+        /**
+         * @brief Write data from host pointer to OpenCL buffer
+         * @param buffer OpenCL buffer to write data to
+         * @param block Whether this call should block and wait until the data operation is finished.
+         * @param offset Byte offset in OpenCL buffer to write data to
+         * @param size nr of bytes to write (e.g. elements*sizeof(datatype))
+         * @param pointerToData Pointer to host memory to read data from. Must be bigger than size.
+         */
         void addWriteBuffer(OpenCLBuffer buffer, bool block, std::size_t offset, std::size_t size, void* pointerToData);
+        /**
+         * @brief Copy data from one OpenCL buffer (source) to another (destination)
+         * @param srcBuffer source OpenCL buffer
+         * @param dstBuffer destination OpenCL buffer
+         * @param srcOffset offset in bytes in source buffer
+         * @param destOffset  offset in bytes in destination buffer
+         * @param size size in bytes (e.g. elements*sizeof(datatype))
+         */
         void addCopyBuffer(OpenCLBuffer srcBuffer, OpenCLBuffer dstBuffer, std::size_t srcOffset, std::size_t destOffset, std::size_t size);
         cl::CommandQueue getHandle() const;
     private:
         cl::CommandQueue m_queue;
 };
 
+/**
+ * @brief Access to OpenCL memory granted to kernel
+ * @ingroup opencl
+ */
 enum class KernelMemoryAccess {
     READ_WRITE = 0,
     READ_ONLY,
     WRITE_ONLY
 };
+/**
+ * @brief Access to OpenCL memory granted to host
+ * @ingroup opencl
+ */
 enum class HostMemoryAccess {
     UNSPECIFIED = 0,
     READ_ONLY,
     WRITE_ONLY,
     //READ_WRITE,
     NONE
 };
+
 /**
  * @brief OpenCL Buffer
  * @ingroup opencl
  */
 class FAST_EXPORT OpenCLBuffer {
     public:
+        /**
+         * @brief Create OpenCL buffer
+         * @param size in bytes (e.g. elements*sizeof(datatype))
+         * @param device device
+         * @param kernelAccess Access to memory to grant to kernel (default READ+WRITE)
+         * @param hostAccess Access to memory to grant to host (default UNSPECIFIED)
+         * @param data Pointer to data on host which is copied to the device (default nullptr/none)
+         */
         OpenCLBuffer(
                 std::size_t size,
                 OpenCLDevice::pointer device,
@@ -257,6 +313,10 @@ class FAST_EXPORT OpenCLBuffer {
                 const void* data = nullptr
         );
         cl::Buffer getHandle() const;
+        /**
+         * @brief Get size of OpenCL buffer in bytes
+         * @return  number of bytes
+         */
         std::size_t getSize() const;
     private:
         cl::Buffer m_buffer;