代码之家  ›  专栏  ›  技术社区  ›  George

从推力到阵列燃料-使用?

  •  -2
  • George  · 技术社区  · 9 年前

    我正在试着更换一些对arrayfire的紧急呼叫,以检查性能。

    我不确定我是否使用了正确的阵列,因为我所取得的结果根本不匹配。

    所以,我使用的推力代码是:

    cudaMalloc( (void**) &devRow, N * sizeof(float) );
    ...//devRow is filled
    
    thrust::device_ptr<float> SlBegin( devRow );
    for ( int i = 0; i < N; i++, SlBegin += PerSlElmts )
    {
        thrust::inclusive_scan( SlBegin, SlBegin + PerSlElmts, SlBegin );
    }
    
    cudaMemcpy( theRow, devRow, N * sizeof(float), cudaMemcpyDeviceToHost );
    //use theRow...
    

    阵列:

    af::array SlBegin( N , devRow );
    for ( int i = 0;i < N; i++,SlBegin += PerSlElmts )
    {
        accum( SlBegin );
    }
    
    cudaMemcpy( theRow, devRow, N * sizeof(float), cudaMemcpyDeviceToHost );
    //use theRow..
    

    我不确定arrayfire如何处理副本: af::array SlBegin( N , devRow ); 。我们有一个设备指针,它从devRow指向SlBegin,但在数组中。。?

    另外,我想问一下使用gfor的问题。 在里面 arrayfire webpage ,它指出

    请勿直接使用此功能;参见GFOR:并行循环。

    然后对于GFOR:

    在当前版本的ArrayFire中禁用了GFOR

    所以,我们不能使用gfor?

    ---------更新---------------------------

    我有一个小的运行示例,显示了不同的结果:

    #include <stdio.h>
    #include <stdlib.h>
    
    #include <cuda.h>
    #include <cuda_runtime.h>
    #include <curand_kernel.h>
    
    #include "arrayfire.h"
    
    #include <thrust/scan.h>
    #include <thrust/host_vector.h>
    #include <thrust/device_vector.h>
    
    __global__ void Kernel( const int N ,float * const devRow )
    {
    
       int i = threadIdx.x;
       if ( i < N )
            devRow[ i ] = i;
    
     }
    
    int main(){
    
        int N = 6;
        int Slices = 2;
        int PerSlElmts = 3;
    
        float * theRow = (float*) malloc ( N * sizeof( float ));
    
        for ( int i = 0; i < N; i ++ )
            theRow[ i ] = 0;
    
        // raw pointer to device memory
        float * devRow;
        cudaMalloc( (void **) &devRow, N * sizeof( float ) );
    
        Kernel<<< 1,N >>>( N , devRow );
        cudaDeviceSynchronize();
    
        // wrap raw pointer with a device_ptr
        thrust::device_ptr<float> SlBegin( devRow );
    
        for ( int i = 0; i < Slices; i++ , SlBegin += PerSlElmts )
            thrust::inclusive_scan( SlBegin, SlBegin + PerSlElmts , SlBegin );
    
        cudaMemcpy( theRow, devRow, N * sizeof(float), cudaMemcpyDeviceToHost );
    
        for ( int i = 0; i < N; i++ )
            printf("\n Thrust accum : %f",theRow[ i ] );
    
    
        //--------------------------------------------------------------------//
        Kernel<<< 1,N >>>( N , devRow );
        cudaDeviceSynchronize();
    
        af::array SlBeginFire( N, devRow );
    
        for ( int i = 0; i < Slices; i++ , SlBeginFire += PerSlElmts )
            af::accum( SlBeginFire );
    
        SlBeginFire.host( theRow );
    
        for ( int i = 0; i < N; i++ )
                printf("\n Arrayfire accum : %f",theRow[ i ] );
    
        cudaFree( devRow );
        free( theRow );
    
    
        return 0;
    
    }
    
    1 回复  |  直到 9 年前
        1
  •  2
  •   shehzan    9 年前

    看起来您试图在2D阵列上运行逐列(ArrayFire中为第0个dim)扫描。下面是一些您可以使用的代码:

    af::array SlBegin(N, devRow);
    af::array result = accum(SlBegin, 0);
    

    以下是示例输出

    A [5 3 1 1]
    0.7402     0.4464     0.7762 
    0.9210     0.6673     0.2948 
    0.0390     0.1099     0.7140 
    0.9690     0.4702     0.3585 
    0.9251     0.5132     0.6814 
    
    accum(A, 0) [5 3 1 1]
    0.7402     0.4464     0.7762 
    1.6612     1.1137     1.0709 
    1.7002     1.2236     1.7850 
    2.6692     1.6938     2.1435 
    3.5943     2.2070     2.8249 
    

    这将对每一列独立运行和包含扫描。

    至于gfor,它已被添加到ArrayFire的开源版本中。由于这个代码库仍然是测试版,因此改进和修复正在迅速进行。所以请关注我们的github页面。