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

来源：互联网发布：苹果5s2g网络怎么改4g 编辑：程序博客网时间：2024/05/19 14:17

简介

在优化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 01 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 02 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 03 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 04 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 05 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

关键点

完成在模版中使用常量的关键在于在.cu文件里的声明，告诉nvcc要编译哪几个常量的函数！

Enjoy！

0 0