找回密码
 立即注册

QQ登录

只需一步,快速开始

查看: 66|回复: 3

Cuda Reduce 示例程序运行问题

[复制链接]
发表于 2018-9-26 15:00:27 | 显示全部楼层 |阅读模式
ESC4000G3
本帖最后由 brfrank 于 2018-9-26 15:04 编辑

正在学习 reduce 思路,从网上下载了一段示例代码:

编译时发现在红色字体 SharedMemory<T> smem; 通不过
Error    63    error: type name is not allowed   

Error    103    error: identifier "smem" is undefined   
Error    14    error: identifier "SharedMemory" is undefined   


请教各位大佬,我这个是怎么回事?


#ifndef __REDUCTION_MAX_KERNEL_H__
#define __REDUCTION_MAX_KERNEL_H__
#include <stdio.h>
#endif // #ifndef __REDUCTION_MAX_KERNEL_H__

template <class T, unsigned int blockSize>
__global__ void reduce_max_kernel(T *g_odata, T *g_idata, unsigned int n)
{

    SharedMemory<T> smem;
    T *sdata = smem.getPointer();

    // perform first level of reduction,

    // reading from global memory, writing to shared memory

    unsigned int tid = threadIdx.x;
   .... 后略


回复

使用道具 举报

发表于 2018-9-26 18:01:15 | 显示全部楼层
Jetson TX2
楼主你好,

(1)shared memory虽然叫shared memory, 但是它是用的__shared__标注的, 而不是Shared Memory.
如果你的SharedMemory是个类的名字之类的, 请确保标识符有效。

(2)不建议在CUDA中使用C++。可能会导致性能上的损失。

(3)我不懂C++。无法提供更多建议。建议你从哪里下载的代码, 就询问哪里的作者。
回复 支持 反对

使用道具 举报

 楼主| 发表于 2018-9-27 09:31:03 | 显示全部楼层
Tesla P100
屠戮人神 发表于 2018-9-26 18:01
楼主你好,

(1)shared memory虽然叫shared memory, 但是它是用的__shared__标注的, 而不是Shared Mem ...

首先感谢版主的回答

这个是Reduce算法的一个常用例子,其来源于
https://devtalk.nvidia.com/defau ... nce/cuda-reduction/

和老外扯英文比较累所以先来中文论坛看看。
其次在NVIDIA CUDA Samples 9.2 以上的 版本也添加了类似
Advanced/reductionMultiBlockCG. Demonstrates single pass reduction using Multi Block Cooperative Groups.
我没安装9.2以上版本,这个例子看不到的,所以想请教一下论坛各位大佬,玩过Reudce算法的看看。

我估计是一个定义问题, SharedMemory<T> 的定义方式
回复 支持 反对

使用道具 举报

 楼主| 发表于 2018-9-27 09:44:55 | 显示全部楼层
本帖最后由 brfrank 于 2018-9-27 09:46 编辑

我在较新的版本的NVIDIA CUDA Samples 找到了完整可执行代码, 确实是个定义问题,本问题结案。感兴趣的童鞋可以参考附件代码

template<class T>
struct SharedMemory
{
    __device__ inline operator       T *()
    {
        extern __shared__ int __smem[];
        return (T *)__smem;
    }

    __device__ inline operator const T *() const
    {
        extern __shared__ int __smem[];
        return (T *)__smem;
    }
};

// specialize for double to avoid unaligned memory
// access compile errors

template<>
struct SharedMemory<double>
{
    __device__ inline operator       double *()
    {
        extern __shared__ double __smem_d[];
        return (double *)__smem_d;
    }

    __device__ inline operator const double *() const
    {
        extern __shared__ double __smem_d[];
        return (double *)__smem_d;
    }
};


/*
    Parallel sum reduction using shared memory
    - takes log(n) steps for n input elements
    - uses n threads
    - only works for power-of-2 arrays
*/

/* This reduction interleaves which threads are active by using the modulo
   operator.  This operator is very expensive on GPUs, and the interleaved
   inactivity means that no whole warps are active, which is also very
   inefficient */

template <class T>
__global__ void
reduce0(T *g_idata, T *g_odata, unsigned int n)
{
    T *sdata = SharedMemory<T>();


// A little disappointed about the domestic technical Forum.



reduction.rar

13.17 KB, 下载次数: 0

回复 支持 反对

使用道具 举报

您需要登录后才可以回帖 登录 | 立即注册

本版积分规则

关闭

站长推荐上一条 /1 下一条

快速回复 返回顶部 返回列表