ReSQL.ru
     
powered by simpleCommunicator - 2.0.59     © 2025 Programmizd 02
Мобильная версия    Контакт    Правила    FAQ    Помощь
Целевая тема:
Создать новую тему:
Автор:
Закрыть
Цитировать
Доб. в избранное | Игнор. тему | Прикреп. тему | Пометить прочит. / непрочит. | Фильтр
Форумы / C++ [игнор отключен] [закрыт для гостей] / агрегация данных на CUDA
6 сообщений из 6, страница 1 из 1
агрегация данных на CUDA
    #38201154
Lepsik
Скрыть профиль Поместить в игнор-лист Сообщения автора в теме
Lepsik
Участник
есть код который умножает каждое значение вертикально стояшего вектора на соответсвуюший елемент в каждой колонке матрицы.

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

Код: 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
27.03.2013, 19:07
| Ответить | Цитировать | Написать  
агрегация данных на CUDA
    #38201475
под какую CC-совместимость
Скрыть профиль Поместить в игнор-лист Сообщения автора в теме
под какую CC-совместимость
Гость
Чтобы сделать быстро, важно под какую 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, 01:43
| Ответить | Цитировать | Написать  
агрегация данных на CUDA
    #38202394
Даже точнее так
Скрыть профиль Поместить в игнор-лист Сообщения автора в теме
Даже точнее так
Гость
Даже точнее так:
Код: 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, 14:46
| Ответить | Цитировать | Написать  
агрегация данных на CUDA
    #38202992
Lepsik
Скрыть профиль Поместить в игнор-лист Сообщения автора в теме
Lepsik
Участник
вылетает с неизвестным кодом ошибки при копировании результат на хост (хотя он размером с матрицу) - возможно что я грид неправильно обьявил?

Код: 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, 18:15
| Ответить | Цитировать | Написать  
агрегация данных на CUDA
    #38203275
Всё просто
Скрыть профиль Поместить в игнор-лист Сообщения автора в теме
Всё просто
Гость
Всё просто:
1. Вы не выделили память под pDevSource, который зачем-то по другому назвали в функции pVerVector. Когда его используете - тогда и ошибка.
2. Ну и не забудьте скопировать исходные данные с хоста на девайс перед вызовом кернел-функции.

А вы сами написали эту GPU-шную функцию?
...
Рейтинг: 0 / 0
28.03.2013, 22:01
| Ответить | Цитировать | Написать  
агрегация данных на CUDA
    #38203402
Lepsik
Скрыть профиль Поместить в игнор-лист Сообщения автора в теме
Lepsik
Участник
Всё простоВсё просто:
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
29.03.2013, 00:02
| Ответить | Цитировать | Написать  
6 сообщений из 6, страница 1 из 1
Форумы / C++ [игнор отключен] [закрыт для гостей] / агрегация данных на CUDA
Найденые пользователи ...
Разблокировать пользователей ...
Читали тему (0):
Читали форум (0):
Пользователи онлайн (0):
x
x
Закрыть


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