@桑薇 这样确实没有问题。
但是我是移植应用时发现这个问题,抽象出的一个小例子。
我想问下如果不改函数的输入和输出,如何正确使用从核加速。
今天我不困 发布的帖子
-
RE: OpenACC制导下的加速代码部分,函数返回值与实际值相差很大,求解释是何原因。
-
RE: OpenACC制导下的加速代码部分,函数返回值与实际值相差很大,求解释是何原因。
@夜深忽梦少年事 test_slave_26.h.c里面确实没有关于my_sum()的定义以及extern声明。
我在中间代码//test_slave.c的my_sum()中加上print,二编执行,发现从核执行时调用的是test_slave.c中的mysum函数。
想请教下,这是为何,以及如何使用openAcc制导,才能在从核代码中得到正确的返回值。 -
RE: OpenACC制导下的加速代码部分,函数返回值与实际值相差很大,求解释是何原因。
//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;
-
RE: OpenACC制导下的加速代码部分,函数返回值与实际值相差很大,求解释是何原因。
@夜深忽梦少年事 感谢小惊喜~
但是从核slave.c中有mysum的定义,而且类型也是正确的。 -
RE: OpenACC制导下的加速代码部分,函数返回值与实际值相差很大,求解释是何原因。
//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); } } }
-
RE: OpenACC制导下的加速代码部分,函数返回值与实际值相差很大,求解释是何原因。
在从核代码中直接print mysum(c,b)值是正确的。赋值到result中,结果就不对了。所以,我怀疑是函数按值返回时出现问题,但是却很不解,看生成中间代码觉得并无不妥。求大神解释下原因
-
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.