максимальное число cuda в блоках и управление распределением

я пишу ядро ​​CUDA, и мне нужно выполнить на этом устройстве:

name: GeForce GTX 480
CUDA capability: 2.0
Total global mem:  1610285056
Total constant Mem:  65536
Shared mem per mp:  49152
Registers per mp:  32768
Threads in warp:  32
Max threads per block:  1024
Max thread dimensions:  (1024, 1024, 64)
Max grid dimensions:  (65535, 65535, 65535)

Ядро в минимальной форме:

_global__ void CUDAvegas( ... )
{
devParameters p;
extern __shared__ double shared[];
int width = Ndim * Nbins;
int ltid = p.lId;
while(ltid < 2* Ndim){
shared[ltid+2*width] = ltid;
ltid += p.lOffset; //offset inside a block
}
__syncthreads();
din2Vec<double> lxp(Ndim, Nbins);

__syncthreads();
for(int i=0; i< Ndim; i++){
  for(int j=0; j< Nbins; j++){
    lxp.v[i][j] = shared[i*Nbins+j];
    }
}
}// end kernel

где Ndim=2, Nbins=128, devParameters — это класс, метод p.lId которого предназначен для подсчета идентификатора локального потока (внутри блока), а din2Cec — это класс для создания вектора dim Ndim*Nbins с помощью новой команды ( в его деструкторе я реализовал соответствующее удаление []). Вывод nvcc:

nvcc -arch=sm_20   --ptxas-options=-v   file.cu -o file.x
ptxas info    : Compiling entry function '_Z9CUDAvegas4LockidiiPdS0_S0_P7sumAccuP17curandStateXORWOWS0_i' for 'sm_20'
ptxas info    : Function properties for  _Z9CUDAvegas4LockidiiPdS0_S0_P7sumAccuP17curandStateXORWOWS0_i
               0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 22 registers, 116 bytes cmem[0], 51200 bytes cmem[2]

количество потоков совместимо с ограничениями мультипроцессоров: максимальное количество разделяемой памяти, максимальное количество регистров на поток и MP и деформации на MP. Если я запускаю 64 потока X 30 блоков (общая память на блок 4128), все в порядке, но если использовать более 30 блоков, я получаю ошибку:

cudaCheckError() failed at file.cu:508 : unspecified launch failure
========= Invalid __global__ read of size 8
=========     at 0x000015d0 in CUDAvegas
=========     by thread (0,0,0) in block (1,0,0)
=========     Address 0x200ffb428 is out of bounds

Я думаю, что это проблема в выделении памяти одного потока, но я не понимаю, каков мой лимит на MP и на общее количество блоков... Кто-нибудь может мне помочь или напомнить правильную тему?

PS: я знаю, что представленное ядро ​​ничего не делает, но это просто для понимания моих проблем с лимитом.


limit memory-management memory cuda register-allocation
person Sim Sca    schedule 23.02.2013    source источник
comment
Если вам нужна помощь в отслеживании выхода за пределы доступной памяти, по крайней мере опубликуйте полный код. Как мы можем узнать, где в вашем ядре генерируется ошибка, без номера строки или кода, который можно было бы скомпилировать и запустить?   -  person talonmies    schedule 23.02.2013

Ответы (1)


arrow_upward
1
arrow_downward

Я думаю, что ошибка, которую вы получаете, является объяснительной. Отмечено, что существует глобальное чтение за пределами памяти типа данных размера 8. Ответственным за чтение за пределами границ является поток (0,0,0) в блоке (1,0,0) . Я подозреваю, что ответственная инструкция - lxp.v[i][j] = shared[i*Nbins+j]; в последнем вложенном цикле for. Вероятно, вы выделяете объем глобальной памяти, который не связан с количеством запускаемых вами блоков, поэтому при запуске слишком большого количества блоков вы получаете такую ​​ошибку.

person Vitality    schedule 24.02.2013
comment
Что касается вашего подозрения по поводу объема памяти, выделяемой одному потоку, я думаю (но, пожалуйста, поправьте меня, если я ошибаюсь), что, пока вы выделяете слишком много регистрового пространства, компилятор начнет сбрасывать память в L1 или глобальную память. - person Vitality; 25.02.2013
comment
Спасибо за ответ, Джек! - person Sim Sca; 28.02.2013