CUDA template kernel 与其他编译器合作编译

最新推荐文章于 2026-06-23 16:49:51 发布

原创最新推荐文章于 2026-06-23 16:49:51 发布 · 1.4k 阅读

0 ·

本内容遵循CC 4.0 BY-SA版权协议

标签

#cuda #c语言 #编译器 #kernel

C/C++ 同时被 3 个专栏收录

63 篇文章

订阅专栏

CUDA

49 篇文章

订阅专栏

技巧

47 篇文章

订阅专栏

本文探讨了在CUDA编程中如何利用C++模板和常量优化Kernel，通过在.cu文件中声明常量函数，使得编译器能够展开循环以提高计算效率。示例代码展示了如何组织源文件，确保模板的正确编译和使用。

简介

在优化Kernel的时候，希望某些变量是常量，例如循环的次数相关的变量。如果次数限制是常量的话，编译器就可以将循环展开。展开的循环，会省掉一些判断，从而节省一些计算时间。

C++的模版中可以使用常量。但是我又不想所有的源代码都由nvcc来编译（其实C++的代码还是调用的host compiler），故此，我写这篇博客来提供一种方法。

代码实例

实例中有三个文件：
main.cpp用host compiler来编译。
cuda_interfaces.cu 用nvcc来编译。

cuda_interfaces.h 是 cuda_interfaces.cu的接口头文件。

// main.cpp
#include <iostream>
#include "cuda_interfaces.h"
int main(int argc, char** argv){
    su::gpu_func<0>();
    su::gpu_func<1>();
    su::gpu_func<2>();
    su::gpu_func<3>();
    su::gpu_func<4>();
    su::gpu_func<5>();
    return EXIT_SUCCESS;
}

// cuda_interfaces.h
#ifndef __CUDA_INTERFACES_H__
#define __CUDA_INTERFACES_H__
namespace su{
    template<int _s> void gpu_func();
}
#endif

// cuda_interfaces.cu
#include <host_defines.h>
#include <device_launch_parameters.h>

#include <iostream>
using namespace std;

namespace su{

    template <int _s>
    __global__ void kernel_func(int *data)
    {
        int x = threadIdx.x + blockIdx.x*blockDim.x;
        if (x < _s){
            data[x] = _s;
        }
        else{
            data[x] = 0;
        }
    }


    template<int _s> void gpu_func()
    {
        int n_threads = 32;
        int *h_data = new int[n_threads];
        int *d_data = NULL;
        cudaMalloc(&d_data, n_threads*sizeof(int));

        kernel_func<_s><<<1, n_threads>>>(d_data);
        cudaMemcpy(h_data, d_data, n_threads*sizeof(int), cudaMemcpyDeviceToHost);

        for (int i = 0; i < n_threads; i++){
            cout << h_data[i] << " ";
        }
        cout << endl;

        // release memory
        delete[] h_data; h_data = NULL;
        cudaFree(d_data); d_data = NULL;
    }

    // note _s only support 0,1,2,3,4,5
    template void gpu_func<0>();
    template void gpu_func<1>();
    template void gpu_func<2>();
    template void gpu_func<3>();
    template void gpu_func<4>();
    template void gpu_func<5>();
}

上述代码执行结果如下：

0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
2 2 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
3 3 3 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
4 4 4 4 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
5 5 5 5 5 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0