小编dar*_*ari的帖子

我在OpenGL程序中使用这个 FXAA Shader进行抗锯齿处理.现在我在CUDA中重新实现了这段代码并进行了测试.生成的图像是相同的,但CUDA版本要慢得多.(Shader使用vsync以60 FPS运行,而CUDA则降至~40 FPS)

这是CUDA代码:

__device__ uchar4 readChar(int x, int y){
    return surf2Dread<uchar4>( surfaceRead, (x)*sizeof(uchar4), (y),cudaBoundaryModeClamp);
}

__device__ uchar4 readFloatBilin2(float x, float y){
    int x1 = floor(x);
    int y1 = floor(y);

    uchar4 z11 = readChar(x1,y1);
    uchar4 z12 = readChar(x1,y1+1);
    uchar4 z21 = readChar(x1+1,y1);
    uchar4 z22 = readChar(x1+1,y1+1);

    float u_ratio = x - x1;
    float v_ratio = y - y1;
    float u_opposite = 1 - u_ratio;
    float v_opposite = 1 - v_ratio;
    uchar4 result = (z11   * u_opposite  + …

我有以下设置:

1. 一个静态库
2. 链接到(1.)的动态库
3. 链接到(1.)和(2.)的可执行文件

来自静态库的代码现在被复制并存在于动态库和可执行文件中.

问题:

Data(全局变量,静态类成员)是否也重复并且可执行文件和dll是否看到相同的数据？

Linux和Windows之间有区别吗？

你怎么解决这个问题？

编辑:

谢谢你的答案,我无法解释我的案件究竟发生了什么.

静态库没有导出/导入标志.动态库帽子导出自己的符号.

视窗:

动态库有一个静态库的文本+数据段的副本.可执行程序无法知道动态库已链接静态库,因为静态库符号中没有可从外部看到.

Linux的:

动态库具有静态库的文本数据段的副本.动态库虽然包含了静态库中自己的符号表中的所有符号(文本和数据). - >可执行文件看到,动态库已经定义了静态库的所有符号,并且没有重新定义它们.

这很糟糕,因为您通常希望在Linux和Windows上具有相同的行为.

1. 共享符号(默认在linux上)

   • 将静态库中的所有符号链接到共享库时,请对其添加dll导出命令. __attribute__ ((dllexport))
   • 将静态库链接到可执行文件时添加dll import命令. __attribute__ ((dllimport))
   • 代码和数据仅驻留在共享库中,并且可以从外部链接

2. Reduntant符号(默认在Windows上)

   • 您需要确保静态库的符号不包含在共享库的符号表中
   • __attribute__ ((visibility ("hidden"))) 在gcc上
   • 链接可执行文件时,无法在任何地方找到静态库中的符号,因此会再次包含它们.

自CUDA 9起,shfl指令已弃用,应由shfl_sync替换.

但是,当他们表现不同时,我该如何更换它们呢？

代码示例:

__global__
static void shflTest(){
    int tid = threadIdx.x;
    float value = tid + 0.1f;
    int* ivalue = reinterpret_cast<int*>(&value);

    //use the integer shfl
    int ix = __shfl(ivalue[0],5,32);
    int iy = __shfl_sync(ivalue[0],5,32);

    float x = reinterpret_cast<float*>(&ix)[0];
    float y = reinterpret_cast<float*>(&iy)[0];

    if(tid == 0){
        printf("shfl tmp %d %d\n",ix,iy);
        printf("shfl final %f %f\n",x,y);
    }
}

int main()
{
    shflTest<<<1,32>>>();
    cudaDeviceSynchronize();
    return 0;
}

输出:

shfl tmp 1084437299 5
shfl final 5.100000 0.000000

我想知道，为什么 NVCC 无法为小矩阵 (N=4) 展开以下 Cholesky 分解内核。

template<typename T, int N>
__device__ inline
void choleskyKernel2(T* C){
#pragma unroll
    for (int i = 0; i < N; i++){
#pragma unroll
        for (int j = 0; j <= i; j++) {
            double s = 0;
#pragma unroll
            for (int k = 0; k < j; k++){
                s += C[i*N+k] * C[j*N+k];
            }
            s = C[i*N+j] - s;
            C[i*N+j] = (i == j) ?
                        sqrt(s) :
                        (1.0 / C[j*N+j] * (s));
        }
    } …

我正在编写一个简单的memcpy内核，以测量GTX 760M的内存带宽并将其与cudaMemcpy（）进行比较。看起来像这样：

template<unsigned int THREADS_PER_BLOCK>
__global__ static
void copy(void* src, void* dest, unsigned int size) {
    using vector_type = int2;
    vector_type* src2 = reinterpret_cast<vector_type*>(src);
    vector_type* dest2 = reinterpret_cast<vector_type*>(dest);

    //This copy kernel is only correct when size%sizeof(vector_type)==0
    auto numElements = size / sizeof(vector_type);

    for(auto id = THREADS_PER_BLOCK * blockIdx.x + threadIdx.x; id < numElements ; id += gridDim.x * THREADS_PER_BLOCK){
        dest2[id] = src2[id];
    }
}

我还计算了达到100％占用率所需的块数，如下所示：

THREADS_PER_BLOCK = 256 
Multi-Processors: 4 
Max Threads per Multi Processor: 2048 
NUM_BLOCKS = 4 * …

我需要解决许多形式为 Ax=0 的小 (n=4) 齐次线性系统，其中 A 是奇异矩阵。我目前正在使用以下代码：

void solve(const matrix_t& A, vector_t& x){
    auto svd = A.jacobiSvd(Eigen::ComputeFullU | Eigen::ComputeFullV);
    auto V = svd.matrixV();
    x = V.col( A.rows() - 1 );
    x.normalize();
}

有没有更快的方法来做到这一点？