Ошибка ядра CUDA при увеличении номера потока

Я разрабатываю ядро ​​​​пересечения лучевой плоскости CUDA.

Предположим, моя структура плоскости (лица):

typedef struct _Face {
    int ID;
    int matID;

    int V1ID;
    int V2ID;
    int V3ID;

    float V1[3];
    float V2[3];
    float V3[3];

    float reflect[3];

    float emmision[3];
    float in[3];
    float out[3];

    int intersects[RAYS];

} Face;

Я вставил всю структуру, чтобы вы могли получить представление о ее размере. RAYS равно 625 в текущей конфигурации. В следующем коде предположим, что размер массива лиц равен, т.е., 1270 (в общем случае - тысяч).

До сих пор я запускал свое ядро ​​очень наивно:

const int tpb = 64; //threads per block
dim3 grid = (n +tpb-1)/tpb; // n - face count in array
dim3 block = tpb;
//.. some memory allocation etc.
theKernel<<<grid,block>>>(dev_ptr, n);

а внутри ядра у меня был цикл:

__global__ void theKernel(Face* faces, int faceCount) {
    int offset = threadIdx.x + blockIdx.x*blockDim.x;
    if(offset >= faceCount)
        return;
    Face f = faces[offset];
    //..some initialization
    int RAY = -1;
    for(float alpha=0.0f; alpha<=PI; alpha+= alpha_step ){ 
        for(float beta=0.0f; beta<=PI; beta+= beta_step ){ 
            RAY++;
            //..calculation per ray in (alpha,beta) direction ...
            faces[offset].intersects[RAY] = ...; //some assignment

Это об этом. Я перебрал все направления и обновил массив faces. Я работал правильно, но вряд ли был быстрее, чем код процессора.

Итак, сегодня я попытался оптимизировать код и запустить ядро ​​с гораздо большим количеством потоков. Вместо 1 нити на грань я хочу 1 нить на луч грани (это означает, что 625 нитей работают на 1 грань). Изменения были простыми:

dim3 grid = (n*RAYS +tpb-1)/tpb;  //before launching . RAYS = 625, n = face count

и само ядро:

__global__ void theKernel(Face *faces, int faceCount){

int threadNum = threadIdx.x + blockIdx.x*blockDim.x;

int offset = threadNum/RAYS; //RAYS is a global #define
int rayNum = threadNum - offset*RAYS;

if(offset >= faceCount || rayNum != 0)
    return;

Face f = faces[offset];
//initialization and the rest.. again ..

И этот код не вообще не работает. Почему? Теоретически должен работать только 1-й поток (из 625 на Face), так почему же это приводит к плохим (почти никаким) вычислениям?

С уважением, эл.


c++ cuda algebra
person emesx    schedule 01.12.2011    source источник

Ответы (2)


arrow_upward
0
arrow_downward

Максимальный размер сетки в любом измерении — 65535 (Руководство по программированию CUDA, Приложение F). Если до изменения размер вашей сетки был 1000, вы увеличили его до 625000. Это больше предела, поэтому ядро ​​не будет работать правильно.

Если вы определяете размер сетки как

dim3 grid((n + tpb - 1) / tpb, RAYS);

тогда все размеры сетки будут меньше предела. Вам также придется изменить способ использования blockIdx в ядре.

person Heatsink    schedule 01.12.2011
comment
Кажется, это оно. К сожалению, прироста производительности не замечено... возможно, из-за того, что внутри ядра остался еще один цикл. Как бы вы разделили (сетку и блоки), имея: 1000 граней, каждая из которых имеет 600 лучей, которые имеют 1000 возможных пересечений? - person emesx; 02.12.2011
comment
Вы уже достигли предела количества потоков, которые могут выполняться параллельно. Добавление большего количества потоков не сделает ядро ​​быстрее. Если вы пытаетесь ускорить работу ядра, может помочь использование другого макета данных для объединения памяти. Объединение памяти описано в руководстве по программированию CUDA. - person Heatsink; 02.12.2011
comment
ну, если у меня есть более 1024 элементов чего-либо, поместить их в блок невозможно, поэтому общая память бесполезна. В моем случае каждое лицо проверяет, какие лица оно может «видеть», кроме самого себя, поэтому часть if(x==self) continue будет различать все потоки; ( - person emesx; 02.12.2011

arrow_upward
0
arrow_downward

Как указал Радиатор, вы, вероятно, превышаете доступные ресурсы. Хорошая идея — проверить после выполнения ядра, не было ли ошибки.

Вот код С++, который я использую:

#include <cutil_inline.h>

void
check_error(const char* str, cudaError_t err_code) {
    if (err_code != ::cudaSuccess)
        std::cerr << str << " -- " << cudaGetErrorString(err_code) << "\n";
}

Затем, когда я задействую ядро:

my_kernel <<<block_grid, thread_grid >>>(args);
check_error("my_kernel", cudaGetLastError());
person Kylo    schedule 01.12.2011