cudaThreadSynchronize()

// 调用CUDA kernel 是非阻塞的，调用kernel语句后面的语句不等待kernel执行完，立即执行。所以在 call_kernel（see kernel.cu）中执行 m5op.dump 是错误的！！！

// REF: https://www.cs.virginia.edu/~csadmin/wiki/index.php/CUDA_Support/Measuring_kernel_runtime

// cudaThreadSynchronize() 暂停调用者的执行，直到前面的 stream operation 执行完毕。

// REF: https://stackoverflow.com/questions/13485018/cudastreamsynchronize-vs-cudadevicesynchronize-vs-cudathreadsynchronize

// C++ thread join 问题，在 kernel.cpp 中也有 join，那么是在 kernel.cpp 中 dump 还是在main.cpp中join后面dump?

// REF: http://en.cppreference.com/w/cpp/algorithm/for_each

// 若 GPU 先执行完毕，在 main.cpp 中join后 dump 似乎合理；若 CPU 先执行完毕，岂不是要阻塞在 cudaThreadSynchronize 处？

// 暂且在 kernel.cp p中 dump！

kernel.cpp

// CPU threads--------------------------------------------------------------------------------------

void run_cpu_threads(T *matrix_out, T *matrix, std::atomic_int *flags, int n, int m, int pad, int n_threads, int ldim, int n_tasks, float alpha

#ifdef CUDA_8_0

    , std::atomic_int *worklist

#endif

    ) {

        std::cout<<"run_cpu_threads start."<<std::endl;

    const int                REGS_CPU = REGS * ldim;

    std::vector<std::thread> cpu_threads;

    for(int i = ; i < n_threads; i++) {

        cpu_threads.push_back(std::thread([=]() {

#ifdef CUDA_8_0

            Partitioner p = partitioner_create(n_tasks, alpha, i, n_threads, worklist);

#else

            Partitioner p = partitioner_create(n_tasks, alpha, i, n_threads);

#endif

            const int matrix_size       = m * (n + pad);

            const int matrix_size_align = (matrix_size + ldim * REGS - ) / (ldim * REGS) * (ldim * REGS);

            for(int my_s = cpu_first(&p); cpu_more(&p); my_s = cpu_next(&p)) {

                // Declare on-chip memory

                T   reg[REGS_CPU];

                int pos      = matrix_size_align -  - (my_s * REGS_CPU);

                int my_s_row = pos / (n + pad);

                int my_x     = pos % (n + pad);

                int pos2     = my_s_row * n + my_x;

// Load in on-chip memory

#pragma unroll

                for(int j = ; j < REGS_CPU; j++) {

                    if(pos2 >=  && my_x < n && pos2 < matrix_size)

                        reg[j] = matrix[pos2];

                    else

                        reg[j] = ;

                    pos--;

                    my_s_row = pos / (n + pad);

                    my_x     = pos % (n + pad);

                    pos2     = my_s_row * n + my_x;

                }

                // Set global synch

                while((&flags[my_s])->load() == ) {

                }

                (&flags[my_s + ])->fetch_add();

                // Store to global memory

                pos = matrix_size_align -  - (my_s * REGS_CPU);

#pragma unroll

                for(int j = ; j < REGS_CPU; j++) {

                    if(pos >=  && pos < matrix_size)

                        matrix_out[pos] = reg[j];

                    pos--;

                }

            }

        }));

    }

    std::for_each(cpu_threads.begin(), cpu_threads.end(), [](std::thread &t) { t.join(); });

    std::cout<<"dump.. after run_cpu_threads end."<<std::endl;

    m5_dump_stats(,);

}

kernel.cu

cudaError_t call_Padding_kernel(int blocks, int threads, int n, int m, int pad, int n_tasks, float alpha,

    T *matrix_out, T *matrix, int *flags

#ifdef CUDA_8_0

    , int l_mem_size, int *worklist

#endif

    ){

        std::cout<<"call_pad start."<<std::endl;

    dim3 dimGrid(blocks);

    dim3 dimBlock(threads);

    Padding_kernel<<<dimGrid, dimBlock

#ifdef CUDA_8_0

        , l_mem_size

#endif

        >>>(n, m, pad, n_tasks, alpha,

        matrix_out, matrix, flags

#ifdef CUDA_8_0

        , worklist

#endif

        );

    cudaError_t err = cudaGetLastError();

    std::cout<<"dump.. after call_pad end."<<std::endl;

    m5_dump_stats(,);

    return err;

}

main.cpp

for(int rep = ; rep < p.n_warmup + p.n_reps; rep++) {

        // Reset

#ifdef CUDA_8_0

        for(int i = ; i < p.n_bins; i++) {

            h_histo[i].store();

        }

#else

        memset(h_histo, , p.n_bins * sizeof(unsigned int));

        cudaStatus = cudaMemcpy(d_histo, h_histo, p.n_bins * sizeof(unsigned int), cudaMemcpyHostToDevice);

        cudaThreadSynchronize();

        CUDA_ERR();

#endif

        std::cout<<"m5 work begin."<<std::endl;

        // Launch GPU threads

        // Kernel launch

        if(p.n_gpu_blocks > ) {

            std::cout<<"launch gpu."<<std::endl;

            cudaStatus = call_Histogram_kernel(p.n_gpu_blocks, p.n_gpu_threads, p.in_size, p.n_bins, n_cpu_bins,

                d_in, (unsigned int*)d_histo, p.n_bins * sizeof(unsigned int));

            CUDA_ERR();

        }

        // Launch CPU threads

        std::cout<<"launch cpu."<<std::endl;

        std::thread main_thread(run_cpu_threads, (unsigned int *)h_histo, h_in, p.in_size, p.n_bins, p.n_threads,

            p.n_gpu_threads, n_cpu_bins);

            std::cout<<"cuda sync."<<std::endl;

        cudaThreadSynchronize();

        std::cout<<"cpu join after cuda sync."<<std::endl;

        main_thread.join();

        //m5_work_end(0, 0);

        std::cout<<"m5 work end."<<std::endl;

    }

巴特西

cudaThreadSynchronize()

最新文章

热门文章