CUDA Lab Course Reference Manual 2011

/Users/sck/cuda/Praktikum_2011/include/lac/texture_matrix.h

Go to the documentation of this file.

#ifndef cublas_Matrix_H
#define cublas_Matrix_H

#include <lac/cublas++.h>
#include <cuda_runtime.h>

// @sect3{Klasse: TextureMatrix}
//
// Diese Klasse verkapselt eine 2D Matrix. Hauptaufgabe ist die Speicherverwaltung auf dem Device.
// Obwohl die Klasse elementweisen Zugriff bietet, sollte dieser nur in Ausnahmefällen benutzt werden.
// Stattdessen sollte die FullMatrix-Klasse als Datenlieferant benutzt werden
struct textureReference;
struct cudaArray;

namespace cublas {

template<typename T>
class TextureMatrix {
    const textureReference * __texref;
    unsigned int __n_rows, __n_cols;
    cudaArray * d_data;
    
    void copy_from(const ::FullMatrix<T> & other);
    void reinit(unsigned int n_rows, unsigned int n_cols);
    void deinit();
    
    void obtain_tex_ref(const char * _texref_name);
    void create_tex_ref();
public:
    TextureMatrix();
    TextureMatrix(unsigned int n_rows, unsigned int n_cols);
    TextureMatrix(::FullMatrix<T> & other);
    
    TextureMatrix(const char * _texref_name);
    TextureMatrix(const char * _texref_name, unsigned int n_rows, unsigned int n_cols);
    TextureMatrix(const char * _texref_name, ::FullMatrix<T> & other);
    
    ~TextureMatrix();

    TextureMatrix<T> & operator = (const ::FullMatrix<T> & other);

    unsigned int n_rows() { return this->__n_rows; };
    unsigned int n_cols() { return this->__n_cols; };
};

}


// @sect3{cublas::TextureMatrix Methoden}
//

// @sect4{Konstruktoren}

template<typename T>
cublas::TextureMatrix<T>::TextureMatrix() : d_data(NULL) {
    create_tex_ref();
    __n_rows = 0;
    __n_cols = 0;
}

template<typename T>
cublas::TextureMatrix<T>::TextureMatrix(unsigned int n_rows, unsigned int n_cols) : d_data(NULL) {
    create_tex_ref();
    reinit(n_rows, n_cols);
}

template<typename T>
cublas::TextureMatrix<T>::TextureMatrix(::FullMatrix<T> & other) : d_data(NULL) {
    create_tex_ref();
    copy_from(other);
}

template<typename T>
cublas::TextureMatrix<T>::TextureMatrix(const char * _texref_name) : d_data(NULL) {
    obtain_tex_ref(_texref_name);
    __n_rows = 0;
    __n_cols = 0;
}

template<typename T>
cublas::TextureMatrix<T>::TextureMatrix(const char * _texref_name, unsigned int n_rows, unsigned int n_cols) : d_data(NULL) {
    obtain_tex_ref(_texref_name);
    reinit(n_rows, n_cols);
}

template<typename T>
cublas::TextureMatrix<T>::TextureMatrix(const char * _texref_name, ::FullMatrix<T> & other) : d_data(NULL) {
    obtain_tex_ref(_texref_name);
    copy_from(other);
}

template<typename T>
cublas::TextureMatrix<T>::~TextureMatrix() {
    if (d_data) {
        deinit();
    }
}

template<typename T>
void cublas::TextureMatrix<T>::reinit(unsigned int _n_rows, unsigned int _n_cols) {
    if (d_data != NULL && (_n_rows != __n_rows || _n_cols != __n_cols)) {
        deinit();
    }
    
    if ((d_data == NULL) && (_n_rows * _n_cols > 0)) {
        __n_rows = _n_rows;
        __n_cols = _n_cols;
        
                // Metadata initialisieren
                cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<T>();
        
                // Initialisiert einen cudaArray, das ist ein 2D array der wegen der Performance optimiert ausgerichtet ist
                cutilSafeCall(cudaMallocArray(&d_data, &channelDesc, __n_cols, __n_rows));

                // Jetzt muss der Array noch an ne Textur gebunden werden
                cutilSafeCall(cudaBindTextureToArray(__texref, d_data, &channelDesc));
        }
}

template<typename T>
void cublas::TextureMatrix<T>::deinit() {
    if (d_data) {
        cutilSafeCall(cudaFreeArray(d_data));
        d_data = NULL;
    }
    __n_rows = 0;
    __n_cols = 0;
}

template<typename T>
void cublas::TextureMatrix<T>::create_tex_ref() {
    
}

template<typename T>
void cublas::TextureMatrix<T>::obtain_tex_ref(const char * _texref_name) {
    cutilSafeCall(cudaGetTextureReference(&__texref, _texref_name));
}

template<typename T>
cublas::TextureMatrix<T> & cublas::TextureMatrix<T>::operator=(const ::FullMatrix<T> & other) {
    copy_from(other);
}

template<typename T>
void cublas::TextureMatrix<T>::copy_from(const ::FullMatrix<T> & other) {
    reinit(other.n_rows(), other.n_cols());
    
    // hier beginnt die Party
    FullMatrixAccessor<T> accessor(other);
    cutilSafeCall(cudaMemcpy2DToArray(this->d_data,                                             /* dst */
                                                                          0, 0,                                                         /* offset */
                                                                          accessor.val(),                                       /* src */
                                                                          this->n_cols() * sizeof(T),           /* src_pitch */
                                                                          this->n_cols() * sizeof(T),           /* width */
                                                                          this->n_rows(),                                       /* height */
                                                                          cudaMemcpyHostToDevice));                     /* direction */
}


#endif

---