CUDA内核访问设备分配结构体无printf输出问题排查

核心问题分析 #

• 野指针访问：主函数仅分配了Network类型的数组内存，但每个Network结构体中的Neurons和Connections指针未指向有效设备内存，内核中直接访问这些指针会触发未定义行为，导致内核崩溃，printf自然无法输出。
• 错误的内存分配方式：cudaMalloc是主机端API，不能在CUDA内核中调用，内核中执行该函数会直接引发错误，导致内核终止。
• 未初始化的结构体成员：Population结构体的num_networks等成员未初始化，虽未直接影响当前逻辑，但属于不良编程习惯。
• 缺乏错误检查：未对CUDA内存操作、内核启动等步骤做错误检查，无法及时定位问题。

解决方案 #

1. 主机端完成所有设备内存分配 #

针对每个Network，在主机端为其Neurons和Connections分配设备内存，再将指针同步到设备端的Network结构体中。

2. 移除内核中的cudaMalloc调用 #

内核仅负责初始化数据，不做内存分配操作。

3. 添加CUDA错误检查 #

在所有CUDA操作后检查错误，快速定位问题。

4. 初始化结构体成员 #

完善Population等结构体的成员初始化。

修改后的完整代码 #

结构体定义（不变） #

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;
};

修正后的内核代码 #

__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("Neuron %d bias: %f\n", i, net->Neurons[i].bias);
    }

    // 初始化输入神经元和连接
    for (int i = 0; i < input_num; ++i) {
        net->Neurons[output_num + i].type = 0; // 输入神经元放在输出神经元之后，避免索引重叠
        net->Neurons[output_num + i].bias = 0.0f;
        net->Neurons[output_num + i].output = 0.0f;
        net->Neurons[output_num + i].input_sum = 0.0f;

        for (int j = 0; j < output_num; ++j) {
            int offset = j + (output_num * i);

            net->Connections[offset].from = output_num + 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("Network %d, Weight [%d]: %f\n", idx, offset, net->Connections[offset].weight);
        }
    }
}

修正后的主函数代码（含错误检查） #

#include <iostream>
#include <curand_kernel.h>
#include <math.h>

// CUDA错误检查宏
#define CHECK_CUDA_ERROR(err) \
    if (err != cudaSuccess) { \
        std::cerr << "CUDA error at " << __FILE__ << ":" << __LINE__ << ": " << cudaGetErrorString(err) << std::endl; \
        exit(EXIT_FAILURE); \
    }

int main() {
    int population_size = 1024;
    int input_num = 10;
    int output_num = 5;
    int neurons_per_net = input_num + output_num;
    int connections_per_net = input_num * output_num;

    // 分配设备端Population结构体
    Population* d_population;
    CHECK_CUDA_ERROR(cudaMalloc(&d_population, sizeof(Population)));

    // 分配设备端Network数组
    Network* d_networks;
    CHECK_CUDA_ERROR(cudaMalloc(&d_networks, sizeof(Network) * population_size));

    // 为每个Network分配Neurons和Connections内存，并同步指针到设备端
    for (int i = 0; i < population_size; ++i) {
        Neuron* d_neurons;
        Connection* d_connections;
        CHECK_CUDA_ERROR(cudaMalloc(&d_neurons, sizeof(Neuron) * neurons_per_net));
        CHECK_CUDA_ERROR(cudaMalloc(&d_connections, sizeof(Connection) * connections_per_net));

        // 更新设备端Network的指针
        CHECK_CUDA_ERROR(cudaMemcpy(&d_networks[i].Neurons, &d_neurons, sizeof(Neuron*), cudaMemcpyHostToDevice));
        CHECK_CUDA_ERROR(cudaMemcpy(&d_networks[i].Connections, &d_connections, sizeof(Connection*), cudaMemcpyHostToDevice));
    }

    // 初始化Population结构体成员并拷贝到设备端
    Population h_population;
    h_population.Networks = d_networks;
    h_population.num_networks = population_size;
    h_population.generation_id = 0;
    CHECK_CUDA_ERROR(cudaMemcpy(d_population, &h_population, sizeof(Population), cudaMemcpyHostToDevice));

    // 启动内核
    int threadsPerBlock = 512;
    int blocks = (population_size + threadsPerBlock - 1) / threadsPerBlock;
    CreateBasePopulation<<<blocks, threadsPerBlock>>>(d_population, population_size, input_num, output_num);
    CHECK_CUDA_ERROR(cudaGetLastError()); // 检查内核启动错误
    CHECK_CUDA_ERROR(cudaDeviceSynchronize()); // 等待内核执行完成，确保printf输出同步

    std::cout << "Population created successfully!" << std::endl;

    // 释放内存：先释放每个Network的Neurons和Connections，再释放Network数组和Population
    for (int i = 0; i < population_size; ++i) {
        Neuron* d_neurons;
        Connection* d_connections;
        CHECK_CUDA_ERROR(cudaMemcpy(&d_neurons, &d_networks[i].Neurons, sizeof(Neuron*), cudaMemcpyDeviceToHost));
        CHECK_CUDA_ERROR(cudaMemcpy(&d_connections, &d_networks[i].Connections, sizeof(Connection*), cudaMemcpyDeviceToHost));
        CHECK_CUDA_ERROR(cudaFree(d_neurons));
        CHECK_CUDA_ERROR(cudaFree(d_connections));
    }
    CHECK_CUDA_ERROR(cudaFree(d_networks));
    CHECK_CUDA_ERROR(cudaFree(d_population));

    return 0;
}

额外说明 #

• 内核中修正了神经元索引重叠问题：原代码中输入、输出神经元都从0开始索引，会覆盖数据，调整后输入神经元放在输出神经元之后（索引从output_num开始）。
• CHECK_CUDA_ERROR宏可快速定位内存分配、内核启动等环节的错误。
• cudaDeviceSynchronize()确保内核所有线程执行完毕，printf的输出会同步到主机端控制台。

内容的提问来源于stack exchange，提问作者ug0x01