Eu tenho um CreateBasePopulation do kernel CUDA onde eu uso printf para imprimir valores de struct dentro do kernel. No entanto, nenhuma saída está sendo impressa quando eu executo o kernel, aqui está o código relevante:
__global__ void CreateBasePopulation(Population* pop, int pop_num, int input_num, int output_num) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx >= pop_num) {
return;
}
Network* net = &pop->Networks[idx];
net->num_neurons = input_num + output_num;
net->num_connections = input_num * output_num;
net->fitness = 0.0f;
curandState state;
curand_init(clock64(), idx, 0, &state);
for (int i = 0; i < output_num; ++i) {
net->Neurons[i].type = 2;
net->Neurons[i].bias = (2.0f * sqrtf((float)input_num) * curand_uniform(&state)) - sqrtf((float)input_num);
net->Neurons[i].output = 0.0f;
net->Neurons[i].input_sum = 0.0f;
printf("%f\n", net->Neurons[i].bias);
}
for (int i = 0; i < input_num; ++i) {
net->Neurons[i].type = 0;
net->Neurons[i].bias = 0.0f;
net->Neurons[i].output = 0.0f;
net->Neurons[i].input_sum = 0.0f;
for (int j = 0; j < output_num; ++j) {
int offset = j + (output_num * i);
net->Connections[offset].from = i;
net->Connections[offset].to = j;
net->Connections[offset].innovationid = offset;
net->Connections[offset].enabled = true;
net->Connections[offset].weight = (2.0f * curand_uniform(&state)) - 1.0f;
printf("Weight [%d]: %f\n", offset, net->Connections[offset].weight);
}
}
}
Eu também tentei alocar memória assim;
....
curandState state;
curand_init(clock64(), idx, 0, &state);
cudaMalloc((void**)&(net->Neurons), sizeof(Neuron) * net->num_neurons);
cudaMalloc((void**)&(net->Connections), sizeof(Connection) * net->num_connections);
....
Mas não dá nenhuma saída (tentei copiar o dispositivo para o host e havia números muito grandes e muito pequenos, então parece que há um problema com o gerenciamento de memória)
Também função principal;
int main() {
int population_size = 1024;
int input_num = 10;
int output_num = 5;
Population* d_population;
cudaMalloc(&d_population, sizeof(Population));
Network* d_networks;
cudaMalloc(&d_networks, sizeof(Network) * population_size);
cudaMemcpy(&(d_population->Networks), &d_networks, sizeof(Network*), cudaMemcpyHostToDevice);
int threadsPerBlock = 512;
int blocks = (population_size + threadsPerBlock - 1) / threadsPerBlock;
CreateBasePopulation<<<blocks, threadsPerBlock>>>(d_population, population_size, input_num, output_num);
cudaDeviceSynchronize();
std::cout << "Population created successfully!" << std::endl;
cudaFree(d_networks);
cudaFree(d_population);
return 0;
}
Também estruturas;
struct Connection {
int innovationid;
int from;
int to;
float weight;
bool enabled;
};
struct Neuron {
int type; //0 = input, 1 = hidden, 2 = output
float input_sum; // Sum of inputs into neuron
float bias;
float output; // Activated output
};
struct Network {
Connection* Connections;
Neuron* Neurons;
int num_connections;
int num_neurons;
float fitness;
};
struct Population {
Network* Networks;
int num_networks;
int generation_id;
};
Primeiro você precisa alocar memória, então a versão inicial do kernel está claramente quebrada.
Usar
cudaMalloc()
no código do dispositivo é um recurso da API de tempo de execução do dispositivo CUDA introduzida para o CUDA Dynamic Parallelism (CDP). Para usar essas APIs com NVCC, é preciso compilar com-rdc=true
o que habilita o código do dispositivo relocável. De acordo com a documentação, também é preciso vincular explicitamente à biblioteca com-lcudadevrt
embora eu não tenha precisado fazer isso para executar o código do OP. Curiosamente, usar outras funções dessa API como devicecudaGetErrorString()
resulta em um erro de compilação (Unresolved extern function '__cudaCDP2GetErrorString'
) na ausência de-rdc=true
. O fato de quecudaMalloc()
se comporta de forma diferente (compila, mas causa um erro de tempo de execução) pode ser um bug.Para evitar essas complicações, pode-se usar
malloc()
/free()
no código do dispositivo (C++new
/delete
também deve funcionar) para obter uma assim chamada alocação de heap de memória do dispositivo. Apenas tenha em mente que a quantidade de memória disponível dessa forma é limitada e que o limite pode ser configurado no host. Essa memória não pode ser usada com nenhuma API de tempo de execução CUDA comocudaMemcpy()
.Como essa memória parece ser subalocada de um buffer em vez de ir para o driver, que é mais flexível devido ao endereçamento virtual, não há como garantir quanto do heap é realmente utilizável, mesmo que se aumente seu tamanho. Então, sem dúvida, não é adequado para o aplicativo do OP com suas muitas alocações (uma por thread). Usar uma thread por bloco para alocar um único buffer para o bloco inteiro pode melhorar a confiabilidade (e o desempenho), mas também aumentar a complexidade do código, especialmente ao liberar a memória mais tarde.
Agora tenho quase certeza de que o mesmo limite de tamanho não se aplica ao
cudaMalloc()
código do dispositivo (não consegui encontrar nenhuma dica de que se aplicaria nos documentos). Portanto, eu esperaria que o dispositivocudaMalloc()
tivesse uma sobrecarga ainda maior do que o dispositivomalloc()
. Ambos os tipos de alocações só podem ser liberados do código do dispositivo, e é por isso que presumi originalmente que o dispositivocudaMalloc()
também usa o heap de memória do dispositivo.Independentemente de qual API você escolher, você deve verificar se a alocação foi bem-sucedida, pelo menos em compilações de depuração (por exemplo, usando
assert()
). Para a API de tempo de execução do dispositivo CUDA, há uma macro de verificação de erros fornecida na parte inferior da resposta superior para "Qual é a maneira canônica de verificar erros usando a API de tempo de execução CUDA?" .