Гость
Войти | Профиль | Очистить
Целевая тема:
Создать новую тему:
Автор:
Действия ...
Форумы / C++ [игнор отключен] [закрыт для гостей] / агрегация данных на CUDA / 6 сообщений из 6, страница 1 из 1
27.03.2013, 19:07
    #38201154
Lepsik
Lepsik
Участник
Скрыть профиль Поместить в игнор-лист Сообщения автора в теме
агрегация данных на CUDA
есть код который умножает каждое значение вертикально стояшего вектора на соответсвуюший елемент в каждой колонке матрицы.

На самом деле мне надо вернуть не результируюшую матрицу, а вектор горизонтальный, содержаший суммы значений каждого столбца уже перемноженной матрицы. Можно это сделать в данном коде?

Код: plaintext
1.
2.
3.
4.
5.
6.
7.
8.
9.
10.
11.
12.
13.
14.
15.
16.
17.
18.
19.
20.
__global__ void AggMtx2Vec( float* pMatrix, float* pResult, const int height, const int width )
 {
    __shared__ float temp[BLOCK_SIZE][BLOCK_SIZE];
 
    int xIndex = blockIdx.x * blockDim.x + threadIdx.x;
    int yIndex = blockIdx.y * blockDim.y + threadIdx.y;
    int tid = threadIdx.x;
    int index  = yIndex * width + xIndex;
 
    if ((xIndex < width) && (yIndex < height))
    {
        temp[threadIdx.x][threadIdx.y] = pMatrix[ index ] * pDevSource[yIndex];
    }
    __syncthreads();
 
    if ((xIndex < height) && (yIndex < width))
    {
        pResult[ index ] = temp[threadIdx.x][threadIdx.y];
    }
 }
...
Рейтинг: 0 / 0
| Ответить | Цитировать | Написать  
28.03.2013, 01:43
    #38201475
под какую CC-совместимость
под какую CC-совместимость
Гость
Скрыть профиль Поместить в игнор-лист Сообщения автора в теме
агрегация данных на CUDA
Чтобы сделать быстро, важно под какую CC-совместимость пишите. В общем случае смотрите исходники scan/scanwarp из CUDPP. Хотя там тоже не оптимально при CC2.0 можно и быстрее.

А если медленно, то как-то так (в pResult будет искомый вектор):
Код: plaintext
1.
2.
3.
4.
5.
6.
7.
8.
9.
10.
11.
12.
13.
14.
15.
16.
17.
18.
19.
20.
21.
22.
23.
24.
25.
26.
27.
28.
29.
__global__ void AggMtx2Vec( float* pMatrix, float* pResult, const int height, const int width )
 {
    __shared__ float temp[BLOCK_SIZE][BLOCK_SIZE];
 
    int xIndex = blockIdx.x * blockDim.x + threadIdx.x;
    int yIndex = blockIdx.y * blockDim.y + threadIdx.y;
    int tid = threadIdx.x;
    int index  = yIndex * width + xIndex;
 
    if ((xIndex < width) && (yIndex < height))
    {
        temp[threadIdx.x][threadIdx.y] = pMatrix[ index ] * pDevSource[yIndex];
    }
    __syncthreads();
 
    if(threadIdx.y == 0) {
      #pragma unroll
      for(int i = 1; i<blockDim.y; ++i) {
          temp[threadIdx.x][0] += temp[threadIdx.x][i];
      }
      __syncthreads();

      if ((xIndex < height) && (yIndex < width))
      {
          pResult[ index ] = temp[threadIdx.x][0];
      }
    }    

 }
...
Рейтинг: 0 / 0
| Ответить | Цитировать | Написать  
28.03.2013, 14:46
    #38202394
Даже точнее так
Даже точнее так
Гость
Скрыть профиль Поместить в игнор-лист Сообщения автора в теме
агрегация данных на CUDA
Даже точнее так:
Код: plaintext
1.
2.
3.
4.
5.
6.
7.
8.
9.
10.
11.
12.
13.
14.
15.
16.
17.
18.
19.
20.
21.
22.
23.
24.
25.
26.
27.
28.
29.
30.
__global__ void AggMtx2Vec( float* pMatrix, float* pResult, const int height, const int width )
 {
    __shared__ float temp[BLOCK_SIZE][BLOCK_SIZE];
 
    int xIndex = blockIdx.x * blockDim.x + threadIdx.x;
    int yIndex = blockIdx.y * blockDim.y + threadIdx.y;
    int tid = threadIdx.x;
    int index  = yIndex * width + xIndex;
 
    if ((xIndex < width) && (yIndex < height))
    {
        temp[threadIdx.x][threadIdx.y] = pMatrix[ index ] * pDevSource[yIndex];
    }
    __syncthreads();
 
    if(threadIdx.y == 0) {
      register float val = 0;
      #pragma unroll
      for(int i = 0; i<blockDim.y; ++i) {
          val += temp[threadIdx.x][i];
      }
      __syncthreads();

      if ((xIndex < height) && (yIndex < width))
      {
          pResult[ index ] = val;
      }
    }    

 }
...
Рейтинг: 0 / 0
| Ответить | Цитировать | Написать  
28.03.2013, 18:15
    #38202992
Lepsik
Lepsik
Участник
Скрыть профиль Поместить в игнор-лист Сообщения автора в теме
агрегация данных на CUDA
вылетает с неизвестным кодом ошибки при копировании результат на хост (хотя он размером с матрицу) - возможно что я грид неправильно обьявил?

Код: plaintext
1.
2.
3.
4.
5.
6.
7.
8.
9.
10.
11.
12.
13.
14.
15.
16.
17.
18.
19.
20.
21.
22.
    const size_t nHight    = 8;
    const size_t nWidth =    8;
    float pHostMatrix[ nHight * nWidth] = {0x00};
    float *pDevMatrix = nullptr;
    err = ::cudaMalloc( (void**)&pDevMatrix, nWidth * sizeof(float) * nHight);

    float pHostResult[nHight * nWidth ] = {0.0f};
    float *pDevResult = nullptr;
    err = ::cudaMalloc((void **)&pDevResult, sizeof(float) * nHight * nWidth);

    dim3 dimGrid(nWidth/TILE_DIM, nHight/TILE_DIM, 1);
    dim3 dimBlock(TILE_DIM, TILE_DIM, 1);

     AggMtx2Vec<<<dimGrid, dimBlock>>>( pDevSource, pDevMatrix, pDevResult, nHight, nWidth);   // calc on GPU

        err = ::cudaMemcpy(pHostResult, pDevResult, sizeof(float) * nHight * nWidth, cudaMemcpyDeviceToHost);

        if (err != cudaSuccess)
        {
            fprintf(stderr, "Failed to copy vector C from device to host (error code %s)!\n", ::cudaGetErrorString(err));  // here is unknow result!!!
            exit(EXIT_FAILURE);
        }



мне кажется я все еше не понимаю принципы адресации

скажем если я просто копирую - это работает:

Код: plaintext
1.
2.
3.
4.
5.
6.
7.
8.
 __global__ void copy( const float* pVerVector, const float* pMatrix, float* pHorVector, const int height, const int width )
 {
    int xIndex = blockIdx.x * blockDim.x + threadIdx.x;
    int yIndex = blockIdx.y * blockDim.y + threadIdx.y;
    int index  = yIndex * width + xIndex;
 
    pHorVector[ index ] = pMatrix[ index ]; // still has the same size as pMatrix
 }



если пытаюсь умножить матрицу на вертикальный вектор - это не работает - та же ошибка при копировании результата на хост
Код: plaintext
1.
2.
3.
4.
5.
6.
7.
8.
 __global__ void copy( const float* pVerVector, const float* pMatrix, float* pHorVector, const int height, const int width )
 {
    int xIndex = blockIdx.x * blockDim.x + threadIdx.x;
    int yIndex = blockIdx.y * blockDim.y + threadIdx.y;
    int index  = yIndex * width + xIndex;
 
    pHorVector[ index ] = pMatrix[ index ] * pVerVector[yIndex]; // still has the same size as pMatrix
 }



source file is attached
...
Рейтинг: 0 / 0
| Ответить | Цитировать | Написать  
28.03.2013, 22:01
    #38203275
Всё просто
Всё просто
Гость
Скрыть профиль Поместить в игнор-лист Сообщения автора в теме
агрегация данных на CUDA
Всё просто:
1. Вы не выделили память под pDevSource, который зачем-то по другому назвали в функции pVerVector. Когда его используете - тогда и ошибка.
2. Ну и не забудьте скопировать исходные данные с хоста на девайс перед вызовом кернел-функции.

А вы сами написали эту GPU-шную функцию?
...
Рейтинг: 0 / 0
| Ответить | Цитировать | Написать  
29.03.2013, 00:02
    #38203402
Lepsik
Lepsik
Участник
Скрыть профиль Поместить в игнор-лист Сообщения автора в теме
агрегация данных на CUDA
Всё простоВсё просто:
1. Вы не выделили память под pDevSource, который зачем-то по другому назвали в функции pVerVector. Когда его используете - тогда и ошибка.
2. Ну и не забудьте скопировать исходные данные с хоста на девайс перед вызовом кернел-функции.

А вы сами написали эту GPU-шную функцию?

1.все там есть в приаттаченом коде
__device__ __constant__ float pDevSource [ 1024 ] = {0x00};

переименовал в последнем варианте для простоты понимания

2. это тоже есть


хотя как оказалось подобная функциональность решается еше и так

Код: plaintext
1.
2.
3.
4.
5.
6.
7.
8.
9.
10.
11.
12.
13.
14.
15.
16.
17.
18.
19.
20.
21.
22.
23.
24.
25.
26.
27.
28.
29.
30.
31.
32.
33.
34.
int KernelCublasSgemv( const float* pVerVector, const  float* pMatrix, float* pHorVector, const size_t hight, const size_t width )
{
    if( nullptr == pVerVector || nullptr == pMatrix || nullptr == pHorVector || 0x00 == hight || 0x00 == width )
    {
        return -1;
    }
    cublasStatus status = ::cublasInit();

    float* pDevMatrix = nullptr;
    
    status = ::cublasAlloc( width * hight, sizeof(float), (void**)&pDevMatrix );

    float* pDevVerVector = nullptr;

    status = ::cublasAlloc( hight, sizeof(float), (void**)&pDevVerVector );

    float* pDevHorVector = nullptr;

    status = ::cublasAlloc( width, sizeof(float), (void**)&pDevHorVector );

    status = ::cublasSetMatrix( hight, width, sizeof(float), pMatrix, width, pDevMatrix, width );

    status = ::cublasSetVector( hight, sizeof(float), pVerVector, 1 , pDevVerVector, 1 );

    status = ::cublasSetVector( width, sizeof(float), pHorVector, 1 , pDevHorVector, 1 );

    ::cublasSgemv('t', width, hight, 1.0, pDevMatrix, width, pDevVerVector, 1, 1.0f, pDevHorVector, 1 );

    status = ::cublasGetVector( width, sizeof(float), pDevHorVector, 1, pHorVector, 1 );

    status = ::cublasGetError();

    return 0;
}
...
Рейтинг: 0 / 0
| Ответить | Цитировать | Написать  
Форумы / C++ [игнор отключен] [закрыт для гостей] / агрегация данных на CUDA / 6 сообщений из 6, страница 1 из 1
Найденые пользователи ...
Разблокировать пользователей ...
Читали тему (0):
Читали форум (0):
Пользователи онлайн (0):
x
x
Закрыть


Просмотр
0 / 0
Close
Debug Console [Select Text]