4
\$\begingroup\$

I'm learning CUDA C++ (and C++). Can you guys give me some feedback on my basic templated matrix implementation?

A brief explanation of my decisions:

  • First things first: I can't use STL classes/functions in my CUDA code.
  • I can't use C++ 14 (or newer) features.
  • I'm using templates to give support to floats, doubles and half precision floats.
  • matrix is meant to be the basic/bare bones implementation; that is a storage for a pointer to a contiguous block of memory, the width and the height of the matrix and the most basic get/set methods.
  • host_matrix is meant to be a "CPU + main memory" implementation. It's just a thin wrapper for matrix plus a constructor/destructor to manage memory.
  • device_matrixis meant to be a "GPU + GPU memory" implementation. Like the host version, it's just a thin wrapper plus some memory management code.
  • The host and device versions don't inherit from matrix would because I can't think of an algorithm that would work properly for a host_matrix and a device_matrix. It would break the substitution principle.

And my most important doubts:

  • Is everything well defined? I'm somewhat concerned that my device_matrix's constructor is working by chance. I'm not sure that the GPU memory layout is compatible with the CPU memory layout and that it's UB to memcpy like I did
  • Does the code make sense?
  • Is the code easy to read/maintain (for C++ standards)? How can I improve the code without making it bloated? By bloated I mean overloading all kinds of operators that won't be used.

#ifndef VK_MATRIX_CUH_INCLUDED
#define VK_MATRIX_CUH_INCLUDED

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <memory>
#include <assert.h>
#include <stddef.h>

namespace vk {

    template<typename T>
    class matrix {
    private:
        T* _values;
        const size_t _width;
        const size_t _height;

    public:
        __host__ __device__ matrix(T* ptr, const size_t width, const size_t height)
            :_values(ptr),
            _width(width),
            _height(height)
        {
            assert(ptr != nullptr);
        }

        __host__ __device__ T* values() const
        {
            return _values;
        }

        __host__ __device__ size_t width() const
        {
            return _width;
        }

        __host__ __device__ size_t height() const
        {
            return _height;
        }

        __host__ __device__ T get(const size_t x, const size_t y) const
        {
            assert(x < _width);
            assert(y < _height);

            size_t offset = y * _width;
            offset += x;
            return _values[offset];
        }

        __host__ __device__ void set(const size_t x, const size_t y, const T value) {
            assert(x < _width);
            assert(y < _height);

            size_t offset = y * _width;
            offset += x;
            _values[offset] = value;
        }
    };

    template<typename T>
    class host_matrix
    {
    private:
        std::shared_ptr<T> _data_pointer;
        std::shared_ptr<matrix<T>> _matrix;

    public:
        host_matrix(const size_t width, const size_t height)
        {
            _data_pointer = std::shared_ptr<T>(new T[width * height], std::default_delete<T[]>());
            _matrix = std::make_shared<matrix<T>>(_data_pointer.get(), width, height);
        }

        T* data() const
        {
            return _data_pointer.get();
        }

        size_t width() const
        {
            return _matrix->width();
        }

        size_t height() const
        {
            return _matrix->height();
        }

        T get(const size_t x, const size_t y) const
        {
            assert(x < width());
            assert(y < height());

            T temp = _matrix->get(x, y);
            return temp;
        }

        void set(const size_t x, const size_t y, const T value)
        {
            assert(x < width());
            assert(y < height());

            _matrix->set(x, y, value);
        }

        T* begin() const
        {
            return _matrix->values;
        }

        T* end() const
        {
            auto offset = _matrix->width * _matrix->height;
            auto last = _matrix->values + offset;
            return last + 1;
        }
    };

    template<typename T>
    class device_matrix
    {
    private:
        T* _data_pointer;
        matrix<T>* _matrix;
        const size_t _width;
        const size_t _height;

    public:
        __host__ device_matrix(const size_t width, const size_t height)
            : _data_pointer(nullptr),
            _matrix(nullptr),
            _width(width),
            _height(height)
        {
            // Allocate gpu memory for the matrix contents
            auto error_code = cudaMalloc(&_data_pointer, width * height * sizeof(T));
            if (error_code != cudaSuccess)
            {
                std::cerr << "Cuda error message: " << cudaGetErrorString(error_code) << std::endl;
                abort();
            }

            // Allocate gpu memory for the struct itself
            error_code = cudaMalloc(&_matrix, sizeof(matrix<T>));
            if (error_code != cudaSuccess)
            {
                std::cerr << "Cuda error message: " << cudaGetErrorString(error_code) << std::endl;
                abort();
            }

            // !!! THIS IS THE PART THAT I'M MORE CONCERNED ABOUT !!! Creates a matrix in host memory and copy it to device
            auto temp = matrix<T>(_data_pointer, width, height);
            error_code = cudaMemcpy(_matrix, &temp, sizeof(temp), ::cudaMemcpyHostToDevice);
            if (error_code != cudaSuccess)
            {
                std::cerr << "Cuda error message: " << cudaGetErrorString(error_code) << std::endl;
                abort();
            }
        }

        __host__ ~device_matrix()
        {
            auto error_code = cudaFree(_data_pointer);
            if (error_code != cudaSuccess)
            {
                std::cerr << "Cuda error message: " << cudaGetErrorString(error_code) << std::endl;
                abort();
            }

            error_code = cudaFree(_matrix);
            if (error_code != cudaSuccess)
            {
                std::cerr << "Cuda error message: " << cudaGetErrorString(error_code) << std::endl;
                abort();
            }
        }

        __host__ __device__ T* data() const
        {
            return _data_pointer;
        }

        __host__ __device__ size_t width() const
        {
            return _width;
        }

        __host__ __device__ size_t height() const
        {
            return _height;
        }

        __device__ T get(const size_t x, const size_t y) const
        {
            assert(x < width());
            assert(y < height());

            T temp = _matrix->get(x, y);
            return temp;
        }

        __device__ void set(const size_t x, const size_t y, const T value)
        {
            assert(x < width());
            assert(y < height());

            _matrix->set(x, y, value);
        }

        __device__ T* begin() const
        {
            return _matrix->values;
        }

        __device__ T* end() const
        {
            auto offset = _matrix->width * _matrix->height;
            auto last = _matrix->values + offset;
            return last + 1;
        }
    };

    template<typename T>
    __host__ void copy_host_to_device(
        const host_matrix<T>& host,
        device_matrix<T>& device)
    {
        assert(host.width() == device.width());
        assert(host.height() == device.height());
        assert(host.data() != nullptr);
        assert(device.data() != nullptr);

        size_t memory_size = host.width() * host.height() * sizeof(T);
        auto error_code = cudaMemcpy(device.data(), host.data(), memory_size, ::cudaMemcpyHostToDevice);
        if (error_code != cudaSuccess)
        {
            std::cerr << "Cuda error message: " << cudaGetErrorString(error_code) << std::endl;
            abort();
        }
    }

    template<typename T>
    __host__ void copy_device_to_host(
        const device_matrix<T>& device,
        host_matrix<T>& host)
    {
        assert(host.width() == device.width());
        assert(host.height() == device.height());
        assert(host.data() != nullptr);
        assert(device.data() != nullptr);

        size_t memory_size = device.width() * device.height() * sizeof(T);
        auto error_code = cudaMemcpy(host.data(), device.data(), memory_size, ::cudaMemcpyDeviceToHost);
        if (error_code != cudaSuccess)
        {
            std::cerr << "Cuda error message: " << cudaGetErrorString(error_code) << std::endl;
            abort();
        }
    }
}

#endif // !VK_MATRIX_CUH_INCLUDED
\$\endgroup\$
2
\$\begingroup\$

class matrix

The class matrix does not own the pointer passed in the constructor. So you should do something to make that clear in the interface so that people don't accidentally pass you dynamically allocated memory.

There is nothing in the standard so I would create a class that represents a raw pointer. Its not

template<typename T>
class raw_ptr
{
    T* ptr
    public:
        explicit raw_ptr(T* ptr): ptr(ptr) {}
        T* operator() const {return ptr;}
};

__host__ __device__
matrix(raw_ptr<T> const& ptr, const size_t width, const size_t height)
        :_values(ptr),
        _width(width),
        _height(height)
    {
        assert(ptr != nullptr);
    }
    // Note:
          T* _values;  // don't need to change this type.
                       // I would just change the interface.

This way you can not just pass a pointer. You need to construct a raw_ptr object that will make people understand that they are supposed to pass an unmanaged pointer across.

   int   data[14] = {};
   matrix<T>  a(raw_ptr<T>(data), 2, 7);

Copying

Because you are not managing the memory you don't technically need to follow the rule of three or five. But this also results in unexpected behavior.

  matrix<int>   b(a);
  a.set(1, 3, 5);
  std::cout << b.get(1, 3) << "\n";   // prints 5 when you would expect 0

As a result I would disable the copy semantics.

  matrix(matrix const&)            = delete;
  matrix& operator=(matrix const&) = delete;

Moving

This should work like normal. There is nothing special needed here.

  matrix(matrix&&)            = default;
  matrix& operator=(matrix&&) = default;

Get and Set

That seems a bit obtuse. When not just one method that returns a reference to the value. You can then use normal assignment to the element.

T& operator()(int x, int y) {
    return values[y * _width + x];
}
T const& operator()(int x, int y) const{
    return values[y * _width + x];
}

int x = a(1,2);    // read a value (get)
a(1,3) = a(1,4);   // write a value from a read value.

With a small amount of work you can get the above to work using [] rather than (). Personally I would do the extra work to get [] working because that is the natural way to use a matrix.

host_matrix

Prefer to use the initializer list

    host_matrix(const size_t width, const size_t height)
    {
        _data_pointer = std::shared_ptr<T>(new T[width * height], std::default_delete<T[]>());
        _matrix = std::make_shared<matrix<T>>(_data_pointer.get(), width, height);
    }

Better if you writ it like this:

    host_matrix(const size_t width, const size_t height)
        : _data_pointer(new T[width * height], std::default_delete<T[]>())
        , _matrix(raw_ptr(_data_pointer.get()), width, height)
    {}

Off by one error:

        auto offset = _matrix->width * _matrix->height;
        auto last = _matrix->values + offset;
        return last + 1;

Think of a 2 * 2 matrix. This is located at [0], [1], [2], [3] so one past the end is at [4] but your function returns [5].

device_matrix

You break the rule of three. You have an owned memory resource. But you don't define the correct copy semantics. Thus you will have multiple calls to cudaFree() when copied objects go out of scope.

You should use a shared pointer and define the destructor just like you did in host_matrix.

class device_matrix { private: std::shared_ptr _data_pointer; std::shared_ptr> _matrix; const size_t _width; const size_t _height;

public:
    __host__ device_matrix(const size_t width, const size_t height)
        : _data_pointer(allocate<T>(width * heigt), [](T* m){cudaFree(m);}),
        _matrix(allocate<matrix<T>(), [](matrix<T>* m){cudaFree(m);}),
        _width(width),
        _height(height)
    {

Use exceptions to kill an app.

The abort() function exits the app. But does not unwind the call stack. I am not sure how the cuda functions affect the GPU if the resources are not freed. But I should not have to worry about that.

If you throw an exception the stack is unwound (thus calling the destructor of all objects). If you have correctly used RAII then these destructors should release your cuda resources. Thus placing your GPU back into a good state.

\$\endgroup\$
  • \$\begingroup\$ Greetings Loki! Thank you A LOT for the feedback. \$\endgroup\$ – Trauer Apr 16 '17 at 14:29

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service, privacy policy and cookie policy

Not the answer you're looking for? Browse other questions tagged or ask your own question.