CUDA C++类中如何利用块与线程实现多级并行？

嘿，作为CUDA新手碰到这类线程协作和共享内存的问题太正常了，我来一步步帮你搞定需求！

核心问题拆解 #

你现在的困境主要来自两个点：

• 类成员不能直接声明为__shared__变量（__shared__是线程块级别的共享内存，不属于单个类实例）
• 原来的kernel写法没有利用块内多线程的协作逻辑，导致多线程执行时重复跑了串行代码部分

1. 正确声明__shared__变量 #

__shared__变量是绑定到线程块的，每个线程块会有自己的独立副本，不能作为类的成员变量存在。你需要把它声明在device函数（比如work）或者kernel内部，这样每个线程块都会分配一块共享内存来存储它。

2. 调整线程协作逻辑 #

要让stuff_that_can_be_done_in_parallel并行执行，同时保证前面的串行代码只跑一次，你需要：

• 让块内只有一个线程（比如threadIdx.x == 0）执行串行计算部分
• 所有线程一起参与并行处理循环
• 用__syncthreads()保证所有线程完成并行计算后，再更新m_data

修改后的完整代码示例 #

#include <cuda_runtime.h>

#define NBLOCKS 16
#define NELEMENTS 1024
#define PARALLEL_SIZE 16

class Test {
public:
    __host__ void set_data(int* d) { 
        // 假设这里是初始化m_data的逻辑
        for (int i = 0; i < NELEMENTS; i++) {
            m_data[i] = d[i];
        }
    }

    __device__ void work(int thread_id) {
        // 声明块级共享内存，每个线程块有自己的shared_data副本
        __shared__ int shared_data[PARALLEL_SIZE];

        while (true) { // 替换成你的实际condition
            int x;
            // 只有块内第一个线程执行串行计算逻辑
            if (thread_id == 0) {
                // ... 这里写你的串行计算代码，最终算出x
                x = blockIdx.x * 10; // 示例值
            }

            // 同步所有线程，确保x已经被计算完成
            __syncthreads();

            // 每个线程处理shared_data的一个元素
            if (thread_id < PARALLEL_SIZE) {
                stuff_that_can_be_done_in_parallel(thread_id, x, shared_data);
            }

            // 等待所有线程完成并行计算
            __syncthreads();

            // 只有一个线程更新m_data（或者根据需求让多个线程参与）
            if (thread_id == 0) {
                // ... 这里写用shared_data更新m_data的逻辑
                for (int i = 0; i < PARALLEL_SIZE; i++) {
                    m_data[i] = shared_data[i];
                }
            }

            // 替换成你的退出条件
            if (blockIdx.x % 2 == 0) break;
        }
    }

    __device__ void stuff_that_can_be_done_in_parallel(int i, int x, int* shared_data) {
        // 示例计算：用x和线程索引填充shared_data
        shared_data[i] = x + i;
    }

    int m_data[NELEMENTS];
    // 移除原来的shared_data成员变量
};

__global__ void kernel(Test* t) {
    // 每个线程调用work，传入自己的线程索引
    t[blockIdx.x].work(threadIdx.x);
}

int main() {
    Test *h_test = new Test[NBLOCKS];
    Test *d_test;

    // 初始化主机端数据
    int dummy_data[NELEMENTS];
    for (int i = 0; i < NELEMENTS; i++) {
        dummy_data[i] = i;
    }
    for (int i = 0; i < NBLOCKS; i++) {
        h_test[i].set_data(dummy_data);
    }

    // 分配设备内存并拷贝数据
    cudaMalloc((void**)&d_test, NBLOCKS * sizeof(Test));
    cudaMemcpy(d_test, h_test, NBLOCKS * sizeof(Test), cudaMemcpyHostToDevice);

    // 启动kernel：NBLOCKS个块，每个块16个线程（对应PARALLEL_SIZE）
    kernel<<<NBLOCKS, PARALLEL_SIZE>>>(d_test);
    cudaDeviceSynchronize();

    // 拷贝结果回主机端
    cudaMemcpy(h_test, d_test, NBLOCKS * sizeof(Test), cudaMemcpyDeviceToHost);

    // 清理资源
    cudaFree(d_test);
    delete[] h_test;

    return 0;
}

关键修改点说明 #

• 移除了类内的shared_data成员，改为在work方法内声明__shared__ int shared_data[PARALLEL_SIZE];，确保每个线程块有独立的共享内存
• work方法新增thread_id参数，让线程知道自己在块内的索引
• 串行计算部分只由thread_id == 0的线程执行，避免重复计算
• 用__syncthreads()在关键节点同步块内所有线程，保证数据一致性
• kernel启动时设置每个块16个线程（对应并行循环的16次迭代），让每个线程处理一个迭代任务

这样就能实现你想要的：每个Test实例的work方法对应一个线程块，块内16个线程并行处理stuff_that_can_be_done_in_parallel部分啦！

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