[cuda] [memcpy3D] Как правильно копировать память?

Ладно... ну нафик. буду использовать обычный cudaMemcpy, без 3д. а с выравниванием памяти и сам справлюсь, вручную.

Можно тред удалять. Хотя как решённый его помечать глупо.

http://sites.google.com/site/deepaknettem/Home/tech-scrapbook/theconceptofpit...

Как я понимаю, у тебя есть 2-х или 3-x мерный массив

ты аллоцируешь его, передавая ссылку на pitch, а индекс элемента считаешь так

idx = y*pitch+x

Не?

Привет. Сейчас вот эта штука снова понадобилась. Только не cudaMalloc3D, а cudaMallocPitch, ну да разница не велика.

Написал простенький код (для теста). И что-то он у меня не работает (заголовки cuda.h и прочие определения опускаю... BLOCK_SIZE = 256). CUDA_CALL - специальный макрос, выводит ошибку.

__global__ void initialize(float *matrix, size_t pitch) {
  uint index = threadIdx.x + blockIdx.x * blockDim.x;
  matrix[pitch*index + index] = index;
}

int main(int argc, char *argv[]) {

  size_t size = 1000;

  float *hostMatrix = new float[size*size];
  float *devMatrix;

  size_t pitch;
  CUDA_CALL(cudaMallocPitch((void**)&devMatrix,&pitch,sizeof(float)*size,size));
  printf("pitch %d\n",pitch);

  dim3 threads(BLOCK_SIZE,1,1);
  dim3 blocks(size/BLOCK_SIZE+1,1,1);
  initialize<<<blocks,threads>>>(devMatrix,pitch);
  CUDA_CALL(cudaGetLastError());
  CUDA_CALL(cudaMemcpy2D(hostMatrix,sizeof(float)*size,devMatrix,pitch,sizeof(float)*size,size,cudaMemcpyDeviceToHost));

  return 0;
}

В общем, ругается на строку cudaMemcpy2D(hostMatrix,sizeof(float)*size,devMatrix,pitch,sizeof(float)*size,size,cudaMemcpyDeviceToHost) - пишет, что:

Error 4: unspecified launch failure

Что я делаю не так?

Методом эксперимента установил, что если задать size поменьше (<=255), то всё работает без ошибок, по крайней мере. (правильность не проверял). Если 256 и более - то вылетает вот с такой ошибкой. Именно копирование памяти.

Помогите пожалуйста, сроки очень горят :( защита скоро... вот...

Кастую в тред Eddy_Em и тех, кто ещё в CUDA соображает...

Хотя.... нет. возможно, проблема не в копировании памяти, а всё-таки в ядре. Закомментил директиву запуска ядра, а копирование оставил - ошибка пропала. Но _ругается_ именно на строку с cudaMemcpy, что странно... может потому что ядро выполняется как бы «в отдельном потоке», а завершается уже после вызова memcpy??..

Подзабыл я уже куду, но сразу бросается в глаза, что двумерный массив копируется в одномерный.

Там, по-видимому, оптимизатор сработал и вообще удалил весь ненужный код ☺

Я, кажется, понял, в чём дело... похоже на ошибку сегментации... Проблема во второй строке в ядре...

Возможно я выхожу за границы массива.

ну я копирую в массив не совсем одномерный, это одномерное представление двумерного же float *hostMatrix = new float[size*size];

Или мне две звёздочки надо писать? Я слышал, такой способ неправильный...

Подзабыл я уже куду

ну мб кто другой подскажет, пусть темка повисит =)

в этом ядре я просто хотел попробовать заполнить главную диагональ матрицы... просто для примера, чтобы понять, как индексы работают... Пока так и не понял.

Если кто укажет на ошибку, буду благодарен...

Приём. который я раньше применял с одномерными массивами... размер блока я делал всегда 256, а размер массива «округлял» вверх до ближайшего кратного 256, то есть если он равен 1000, я добавлял до 1024... ядро работало (параллельно) на 1024 элементах, а затем эти «лишние» 24 элемента я просто отбрасывал (не копировал в host-память)

Хотелось бы аналогично поступить в случае двумерного массива... так проще работать...

Зачем 2D? Просто cudaMemcpy…

Кстати, а зачем у тебя используются макросы CUDA_CALL?

На вскидку dinn, max_udoff, fads, CrossFire

Эээ. А затем, что если я правильно понял, как работает это «выравнивание» - оно добавляет в каждой строке «лишние элементы»... Вот.

Да и проблема, как оказалось, не в копировании. А всё-таки в ядре... обращаюсь за границы массива.

В принципе обращаться за границы массива для меня по идее не страшно... только для этого надо выделить памяти побольше... чтобы за границы не выходить.

Или мне просто надо мыслить по-другому =)

Не мучайся: возьми готовый пример из их dev-kit'а. И не усложняй.

размер блока я делал всегда 256

Лучше при запуске опросить видеокарту (узнать основные параметры) и установить оптимальные размеры блоков и количество потоков - чтобы ядра не простаивали.

#define CUDA_CALL(x) do { if ((x) != cudaSuccess ) {\
  printf("Error at %s:%d\n%d: %s\n",__FILE__,__LINE__,x,cudaGetErrorString(x));\
  exit(EXIT_FAILURE);}} while(0)

Скопипастил, емнип, из мануала по curand. только немного переделал под себя.

В принципе обращаться за границы массива для меня по идее не страшно...

А всякие полосочки, ползущие по монитору, не напрягают? А иногда вообще видеокарта «отваливается»…

Лучше при запуске опросить видеокарту (узнать основные параметры) и установить оптимальные размеры блоков и количество потоков - чтобы ядра не простаивали.

Мне сейчас нужно хотя бы чтобы просто работало =) для моей видеокарты 256 и есть оптимальный размер... если я правильно подсчитывал...

для других мне не надо пока... всё равно их под рукой нет

А, понятно. Просто оно похоже на макрос из кудовских примеров - думал, может ты в них и начал собирать (что не очень хорошо, т.к. без этих самых примеров собираться не будет).

«полосочек» не видел... экран мерцает при запуске, да.. Не напрягает, нет.

Если выделить заранее больше памяти, то тогда никакой блок не будет выходить за границы... Наверное, cudaMallocPitch этой проблемы не решает, и надо самому подсчитывать, округлять до 256^N... этим и займусь

для других мне не надо пока... всё равно их под рукой нет

Ну, у меня дома и на работе разные карты - поэтому сразу и делал так.

А с массивами советую-таки не париться и сделать все проще.

Кстати, раз там двумерная матрица, то и blocks/threads стоило бы сделать двумерными, а в ядре вычислять координату по x/y переменным. Все в примерах есть, да и в мануале разжевано.

«полосочек» не видел...

А у меня поначалу такое часто было: ядро-то сегфолт не выдает, а тупо пишет данные куда-то черт-те куда. И в конце-концов видеокарточка запросто может послать тебя подальше.

Но _ругается_ именно на строку с cudaMemcpy, что странно...

Поскольку у ядра нет возвращаемого значения, код ошибки оставляется на потом, до следующего вызова API. Его можно вытащить с помощью ф-ции cudaGetLastError(). У меня есть такой макрос:

#define CHECKKERNEL { cudaError __ret = cudaGetLastError(); if(__ret!=cudaSuccess) {formatCudaError(__FILE__, __LINE__, __ret, ""); return __ret; } }

// Example:
someKernel<<<gridDim, blockDim>>>(...); CHECKKERNEL;

formatCudaError - просто функция, которая красиво форматирует отчёт об ошибке, делается по вкусу.

cast perestoronin

Да я вроде так и делал (см. выше). Не знаю, почему не сработало... Не столь уж важно теперь... Ошибку локализовал всё равно, осталось придумать, как избавиться от неё :)

initialize<<<blocks,threads>>>(devMatrix,pitch);
CUDA_CALL(cudaGetLastError());

попробую. Помню, было пробовал работать с трёхмерными блоками, и намучился я что-то с ними, не работало ни хрена =)

переделал всё в одномерные (я вообще целиком весь алгоритм там переделал, долго рассказывать), в общем, всё работает и считает как надо. Прирост в 19 раз по сравнению с CPU (правда, задействовано одно ядро из четырёх только), ну это профит (для ВКР вполне себе, это не для работы ;) )

А это вообще другая задача... тут с матрицами хочу поэкспериментировать.

Сдаётся мне, что при том, как ты считаешь размер грида, index может быть больше size.

__global__ void initialize(float *matrix, size_t pitch) {
  uint index = threadIdx.x + blockIdx.x * blockDim.x;
  matrix[pitch*index + index] = index;
}

Можно попробовать поставить

if(index>=size) return;

Помню, было пробовал работать с трёхмерными блоками, и намучился я что-то с ними, не работало ни хрена =)

У тебя, наверное, что-нибудь вроде моей домашней видюшки (9600) - мегакастрированная CUDA получается.

Прирост в 19 раз по сравнению с CPU (правда, задействовано одно ядро из четырёх только)

У меня разные алгоритмы давали прирост где-то в районе одного порядка (на этой видеокарте). Но на CPU параллелил (часть - вручную, часть - при помощи openmp). Кстати, советую openmp использовать - удобная надстройка над pthreads. Сам вот все иногда вспоминаю про свою поделку, но никак не найду достаточно времени (там только чтобы все вспомнить пару дней убить надо). А хотелось бы не только понадобавлять кучу всяких алгоритмов, но и автоматизировать выбор CPU/GPU (если не хватает оперативки видеокарты, то переключать выполнение кода на CPU).

Ну вот я попробовал использовать всё-таки двумерные блоки. Простенький код, который должен заполнить матрицу нулями.

__global__ void initialize(float *matrix, size_t pitch) {
  uint ix = threadIdx.x + blockIdx.x * blockDim.x;
  uint iy = threadIdx.y + blockIdx.y * blockDim.y;
  matrix[iy*pitch + ix] = 0;
}

int main(int argc, char *argv[]) {

  size_t size = 1000;
  size_t realsize = 16 * (size/16 + 1);

  float *hostMatrix = new float[size*realsize];
  float *devMatrix;

  size_t pitch;
  CUDA_CALL(cudaMallocPitch((void**)&devMatrix,&pitch,sizeof(float)*size,size));

  dim3 threads(16,16,1);
  dim3 blocks(realsize/16,realsize/16,1);
  initialize<<<blocks,threads>>>(devMatrix,pitch);
 CUDA_CALL(cudaMemcpy2D(hostMatrix,sizeof(float)*size,devMatrix,pitch,sizeof(float)*size,size,cudaMemcpyDeviceToHost));

  return 0;
}

Отвечает ошибкой 4: unspecified launch failure

ага, 9800 gt, правильно. CUDA у меня только в академических целях, поэтому пока не горю желанием обновлять ;) может, потом куплю что-то помощнее.... А сейчас денег жалко.

Кстати, советую openmp использовать

спс, знаю =) этому нас на параллельном программировании во 2 курсе универа учили... её тоже пользовал. сейчас просто нет цели «побыстрее», а есть цель, чтобы на видяхе получалось быстрее, чем на CPU... или хотя бы просто чтобы было два алгоритма, и они работали :)

openmp действительно удобная штука и очень просто циклы for особенно параллелить. profit...

Ну да, правильно... может быть и больше...

Попробовал добавить

if ((ix>=size) || (iy>=size)) {
    return;
  }

В ядре адрес вычисляется явно как-то нехорошо...

ААА!!! Типун вам на язык :( артефакты на экране появились... вот только что. До этого полтора года на куде писал понемногу и ничего не появлялось. Грусть, тоска, печаль :(

вроде по мануалу... сейчас ещё раз перечитаю и перепроверю.. я просто переписал через квадратные скобки... в теории одно и тоже. сейчас попробую тупо скопипастить какой-нить пример.

Говорил же: неправильно ты, Дядя Федор, колбасу ешь смещение вычисляешь!

Прикольные такие белые точечьки аки снежинки на экране появились... пруф - http://dl.dropbox.com/u/17976101/snap4.png

Неправильное выделение памяти или копирование. Такой код работает без сегфолта:

#include <iostream>
#include <cuda.h>

#define size 1000

__global__ void initialize(float *matrix, size_t pitch) {
  uint index = threadIdx.x + blockIdx.x * blockDim.x;
  if(index<size)
  matrix[size*index + index] = 0.0f;
}

int main(int argc, char *argv[]) {

  int BLOCK_SIZE = 256;

  float *hostMatrix = new float[size*size];
  float *devMatrix;

  size_t pitch;
  //(cudaMallocPitch((void**)&devMatrix,&pitch,sizeof(float)*size,size));
  cudaMalloc(&devMatrix, sizeof(float)*size*size);
  if(cudaGetLastError()!=cudaSuccess)
      std::cout<<"FAIL!";

  dim3 threads(BLOCK_SIZE,1,1);
  dim3 blocks(size/BLOCK_SIZE+1,1,1);
  initialize<<<blocks,threads>>>(devMatrix,pitch);
  if(cudaGetLastError()!=cudaSuccess)
      std::cout<<"FAIL!";
  (cudaMemcpy2D(hostMatrix,sizeof(float)*size,devMatrix,pitch,sizeof(float)*size,size,cudaMemcpyDeviceToHost));
  //cudaMemcpy(hostMatrix, devMatrix, sizeof(float)*size*size, cudaMemcpyDeviceToHost);
  if(cudaGetLastError()!=cudaSuccess)
      std::cout<<"FAIL!";

  return 0;
}

/me bad

Такой вот работает без сегфолта:

#include <iostream>
#include <cuda.h>

#define size 1000

__global__ void initialize(float *matrix, size_t pitch) {
  uint index = threadIdx.x + blockIdx.x * blockDim.x;
  if(index<size)
  matrix[size*index + index] = 0.0f;
}

int main(int argc, char *argv[]) {

  int BLOCK_SIZE = 256;

  float *hostMatrix = new float[size*size];
  float *devMatrix;

  size_t pitch;
  //(cudaMallocPitch((void**)&devMatrix,&pitch,sizeof(float)*size,size));
  cudaMalloc(&devMatrix, sizeof(float)*size*size);
  if(cudaGetLastError()!=cudaSuccess)
      std::cout<<"FAIL!";

  dim3 threads(BLOCK_SIZE,1,1);
  dim3 blocks(size/BLOCK_SIZE+1,1,1);
  initialize<<<blocks,threads>>>(devMatrix,pitch);
  if(cudaGetLastError()!=cudaSuccess)
      std::cout<<"FAIL!";
  //(cudaMemcpy2D(hostMatrix,sizeof(float)*size,devMatrix,pitch,sizeof(float)*size,size,cudaMemcpyDeviceToHost));
  cudaMemcpy(hostMatrix, devMatrix, sizeof(float)*size*size, cudaMemcpyDeviceToHost);
  if(cudaGetLastError()!=cudaSuccess)
      std::cout<<"FAIL!";

  return 0;
}

Будем считать, что решено...

Скопипастил (почти из примера) -

float* pElement = (float*)((char*)matrix + iy * pitch) + ix

Работает без сегфолта, да... Не совсем понятно, зачем тут приведение типов указателей, но, видимо, это ключевое... А вообще начал понимать... ведь размер char - это один байт, а float - это четыре байта!...

В общем, буду впредь аккуратнее относиться к «ненужному» коду :)

ок, спасибо, почитаю и сравню =)

всем спасибо, извиняюсь за беспокойство, да и время уже позднее, всем снофф.

а мне предстоит бессонная ночка с мануалами по cuda и линейной алгеброй %(

Не совсем понятно, зачем тут приведение типов указателей

Потому что pitch имеет тип size_t, а не количество элементов массива!

Шаббат же! Побойся Патрега! Выпей пивка, расслабься. А с утра с новыми силами return to the friday…

Ну это ты погорячился, я к разработке на CUDA никакого отношения не имею :)

пардон, ошибся значит. если ещё кто тут лишний, не обижайтесь =)

Ух ты, темка всплыла. Пардон, что не помог, я тут редко весьма.

Я вроде разобрался. и вроде бы тоже написал код, который обращает матрицу....

правда пока не всё идеально работает, но тут уже проблемы в другом.

Самое главное, алгоритм на gpu работает действительно быстрее =)

вот проверил для матрицы 1000x1000:

cpuTime 13961.235

gpuTime 6531.757

В миллисекундах. С учётом того, что алгоритм корявый и gpu постоянно обменивается с host-памятью %) в силу того, что так требуется... попробую от этого избавиться.