OpenACC制导下的加速代码部分,函数返回值与实际值相差很大,求解释是何原因。



  • 测试代码如下:
    //test.c
    #include<stdio.h>
    #include<stdlib.h>
    #include<math.h>

    double mysum(double f,double m)
    {
    return f + m;
    }

    int main(){
    int i;
    double result[10] = {0};
    #pragma acc parallel loop copy(result) private(i)
    for (i = 0 ; i < 10; i++){
    double c = 0.577577;
    double b = 36.960000;
    result[i] = c+b;
    }
    for (i = 0 ; i < 10; i++){
    printf("%lf\n",result[i]);
    }

    #pragma acc parallel loop copy(result) private(i)
    for (i = 0 ; i < 10; i++){
    double c = 0.577577;
    double b = 36.960000;
    result[i] = mysum(c,b);
    }
    for (i = 0 ; i < 10; i++){
    printf("%lf\n",result[i]);
    }
    return 0;
    }

    swacc test.c
    $ bsub -I -q q_sw_expr -N 1 ./a.out
    Job <43049475> has been submitted to queue <q_sw_expr>
    waiting for dispatch ...
    dispatching ...
    37.537577
    37.537577
    37.537577
    37.537577
    37.537577
    37.537577
    37.537577
    37.537577
    37.537577
    37.537577
    1387858552.000000
    1387858552.000000
    1387858552.000000
    1387858552.000000
    1387858552.000000
    1387858552.000000
    1387858552.000000
    1387858552.000000
    1387858552.000000
    1387858552.000000
    Job 43049475 has been finished.



  • 在从核代码中直接print mysum(c,b)值是正确的。赋值到result中,结果就不对了。所以,我怀疑是函数按值返回时出现问题,但是却很不解,看生成中间代码觉得并无不妥。求大神解释下原因



  • 十有八九是从核没能成功识别函数的返回值类型.
    编译参数加个-keep看一下中间文件是不是看不到mysum的定义.
    那么默认情况下c编译器会认为返回值和参数都是int.
    然后就凉了...
    另外, 贴代码的时候在一堆代码的前面加```c, 后面加```有惊喜.
    最好把中间代码贴上来.



  • //test_host.c

    #include <swacc_host.h> 
    #include <swacc_host.h> 
    #include<stdio.h>
    #include<stdlib.h>
    #include<math.h>
    
    inline double mysum(double f,double m)
    {
      return f + m;
    }
    void slave_test_13(void *argv);
    void slave_test_26(void *argv);
    
    int main()
    {
      int i;
      double result[10] = {(0)};
      acc_set_num_gangs(-1,"test.c",13);
      acc_set_num_workers(-1,"test.c",13);
    {
        acc_c2rtl_t _ptr;
        struct _swacc_pointer {
        double (*_extern_result)[10];}_pointer;
        _pointer . _extern_result = &result;
        void (*func)();
        void *data;
        int num_threads;
        int priv_size;
        int ldm_size;
        _ptr . func = &slave_test_13;
        _ptr . data = &_pointer;
        _ptr . priv_size = 4;
        _ptr . ldm_size = 0;
        _ptr . num_threads = acc_get_num_spe();
        ACC_parallel_start(_ptr);
        ACC_join();
      }
      for (i = 0; i < 10; i++) {
        printf("%lf\n",result[i]);
      }
      acc_set_num_gangs(-1,"test.c",26);
      acc_set_num_workers(-1,"test.c",26);
    {
        acc_c2rtl_t _ptr;
        struct _swacc_pointer {
        double (*_extern_result)[10];}_pointer;
        _pointer . _extern_result = &result;
        void (*func)();
        void *data;
        int num_threads;
        int priv_size;
        int ldm_size;
        _ptr . func = &slave_test_26;
        _ptr . data = &_pointer;
        _ptr . priv_size = 4;
        _ptr . ldm_size = 0;
        _ptr . num_threads = acc_get_num_spe();
        ACC_parallel_start(_ptr);
        ACC_join();
      }
      for (i = 0; i < 10; i++) {
        printf("%lf\n",result[i]);
      }
      return 0;
    }
    
    

    //test_slave.c 从核函数返回类型仍为double

    #include<math.h>
    #include<stdlib.h>
    #include<stdio.h>
    #include <swacc_slave.h> 
    double mysum(double f,double m);
    /* REQUIRED CPP DIRECTIVES */
    /* REQUIRED DEPENDENT DECLARATIONS */
    /* OUTLINED FUNCTION */
    extern __thread_local_fix void *acc_slave_ustack_ptr;
    extern __thread_local_fix void *acc_slave_ustack_ptr_bak;
    
    double mysum(double f,double m)
    {
      return f + m;
    }
    
    

    //test_slave_26.c

    #include <math.h> 
    #include <sys/time.h> 
    #include <swacc_slave.h> 
    #include </home/export/online1/nsccwuxi_ict/chenl/xuy/ablation/sw/test/test_slave_26.h.c> 
    
    void slave_test_26(void *argv)
    {
      int acc_cp_0_size;
      int acc_cp_1_size;
      acc_check_ldm(0 + sizeof(c) + sizeof(b) + sizeof(_ldm_result) + sizeof(acc_data_ret));
      acc_sync_pe_m2l_nostride_copy(&_swacc_ldm_data0x1603560 . _pointer,argv,8);
      acc_myid = acc_get_thread_num();
      acc_corenum = acc_get_num_threads();
      acc_cp_1_size = 1 * 8;
      acc_cp_0_size = 8;
    {
        for (acc_blockindex_i0_0 = acc_myid * 1; acc_blockindex_i0_0 <= 9; acc_blockindex_i0_0 += 1 * acc_corenum) {
          _swacc_ret = acc_sync_pe_m2l_nostride_copy(_ldm_result,&result[acc_blockindex_i0_0],acc_cp_0_size);
          for (acc_index = 0; acc_index <= 0; acc_index += 1) {
            c = 0.577577;
            b = 36.960000;
            _ldm_result[acc_index] = mysum(c,b);
          }
          _swacc_ret = acc_sync_pe_l2m_nostride_copy(&result[acc_blockindex_i0_0],_ldm_result,acc_cp_1_size);
        }
      }
    }
    


  • @夜深忽梦少年事 感谢小惊喜~
    但是从核slave.c中有mysum的定义,而且类型也是正确的。



  • @今天我不困 /home/export/online1/nsccwuxi_ict/chenl/xuy/ablation/sw/test/test_slave_26.h.c里面有my_sum的定义吗? 是要调用函数的文件里面以某种方式找得到这个定义(extern double或者在某个头文件里面)...



  • //test_slave_26.h.c

    struct _swacc_pointer
    {
      double (*_extern_result)[10];
    }
    ;
    typedef struct _swacc_private {
    int acc_p_i;}*__z;
    typedef struct _swacc_ldm {
    struct _swacc_pointer _pointer;
    double c;
    double b;
    int acc_index;
    int acc_lower;
    int acc_upper;
    int acc_myid;
    int acc_corenum;
    double _ldm_result[1];
    int acc_blockindex_i0_0;
    int _swacc_ret;
    int acc_data_ret;}__t;
    extern __thread_local_fix void *accr_priv_ptr;
    extern __thread_local_fix void *accr_ldm_ptr;
    #define result  (*((_swacc_ldm_data0x1603560)._pointer._extern_result))
    #define acc_p_i (((__z)accr_priv_ptr)->acc_p_i)
    #define acc_data_ret    ((_swacc_ldm_data0x1603560).acc_data_ret)
    #define _swacc_ret  ((_swacc_ldm_data0x1603560)._swacc_ret)
    #define acc_blockindex_i0_0 ((_swacc_ldm_data0x1603560).acc_blockindex_i0_0)
    #define _ldm_result ((_swacc_ldm_data0x1603560)._ldm_result)
    #define acc_corenum ((_swacc_ldm_data0x1603560).acc_corenum)
    #define acc_myid    ((_swacc_ldm_data0x1603560).acc_myid)
    #define acc_upper   ((_swacc_ldm_data0x1603560).acc_upper)
    #define acc_lower   ((_swacc_ldm_data0x1603560).acc_lower)
    #define acc_index   ((_swacc_ldm_data0x1603560).acc_index)
    #define b   ((_swacc_ldm_data0x1603560).b)
    #define c   ((_swacc_ldm_data0x1603560).c)
    __thread_kernel("test_slave_26.h")__t _swacc_ldm_data0x1603560;
    
    


  • @夜深忽梦少年事 test_slave_26.h.c里面确实没有关于my_sum()的定义以及extern声明。
    我在中间代码//test_slave.c的my_sum()中加上print,二编执行,发现从核执行时调用的是test_slave.c中的mysum函数。
    想请教下,这是为何,以及如何使用openAcc制导,才能在从核代码中得到正确的返回值。



  • This post is deleted!


  • 试试

     void mysum(double f, double m, double *result)
    {
       *result  =  f + m;
    }
    
    

Log in to reply