summaryrefslogtreecommitdiff
path: root/modules/remote-hip/examples/hip_transfer_data.cpp
diff options
context:
space:
mode:
Diffstat (limited to 'modules/remote-hip/examples/hip_transfer_data.cpp')
-rw-r--r--modules/remote-hip/examples/hip_transfer_data.cpp75
1 files changed, 54 insertions, 21 deletions
diff --git a/modules/remote-hip/examples/hip_transfer_data.cpp b/modules/remote-hip/examples/hip_transfer_data.cpp
index 18c82df..3f02d87 100644
--- a/modules/remote-hip/examples/hip_transfer_data.cpp
+++ b/modules/remote-hip/examples/hip_transfer_data.cpp
@@ -1,20 +1,22 @@
+#include <forstio/codec/data_raw.hpp>
+
#include "../c++/remote.hpp"
#include "../c++/transfer.hpp"
#include <iostream>
-__global__ void print_value(saw::data<saw::schema::Int16,saw::encode::Native>* val){
+__global__ void print_value(saw::data<saw::schema::Int16,saw::encode::NativeRaw>* val){
int v = val->get();
printf("Hello world: %d\n", v);
}
-__global__ void print_array_vals(saw::data<saw::schema::Array<saw::schema::Int16>* val){
+__global__ void print_array_vals(saw::data<saw::schema::Array<saw::schema::Int16>, saw::encode::NativeRaw>* val){
uint64_t orig_len = val->size();
long len = (long) orig_len;
printf("Array size: %ld\n", len);
- for(uint64_t i = 0; i < orig_len; +i){
- int v = val->at(i);
+ for(uint64_t i = 0; i < orig_len; ++i){
+ int v = val->at(i).get();
printf("%d ", v);
}
printf("\n");
@@ -35,28 +37,59 @@ saw::error_or<void> real_main(){
}
auto& addr = eo_addr.get_value();
- auto eo_dat_srv = rmt.data_listen<sch::Int16, encode::Native>(*addr);
- if(eo_dat_srv.is_error()){
- return std::move(eo_dat_srv.get_error());
- }
- auto& dat_srv = eo_dat_srv.get_value();
+ {
+ auto eo_dat_srv = rmt.data_listen<sch::Int16, encode::NativeRaw>(*addr);
+ if(eo_dat_srv.is_error()){
+ return std::move(eo_dat_srv.get_error());
+ }
+ auto& dat_srv = eo_dat_srv.get_value();
- data<sch::Int16> val{42};
- id<sch::Int16> id_val{0u};
- auto eo_send = dat_srv->send(val, id_val);
- if(eo_send.is_error()){
- return std::move(eo_send.get_error());
- }
+ data<sch::Int16,encode::NativeRaw> val{42};
+ id<sch::Int16> id_val{0u};
+ auto eo_send = dat_srv->send(val, id_val);
+ if(eo_send.is_error()){
+ return std::move(eo_send.get_error());
+ }
+
+ auto eo_dfind = dat_srv->find(id_val);
+ if(eo_dfind.is_error()){
+ return std::move(eo_dfind.get_error());
+ }
+ auto dfind = eo_dfind.get_value();
- auto eo_dfind = dat_srv->find(id_val);
- if(eo_dfind.is_error()){
- return std::move(eo_dfind.get_error());
+ auto& v = dfind();
+
+ print_value<<<dim3(2),dim3(2),0,hipStreamDefault>>>(*(v.get_device_data()));
}
- auto dfind = eo_dfind.get_value();
- auto& v = dfind();
+ {
+ auto eo_dat_srv = rmt.data_listen<sch::Array<sch::Int16>, encode::NativeRaw>(*addr);
+ if(eo_dat_srv.is_error()){
+ return std::move(eo_dat_srv.get_error());
+ }
+ auto& dat_srv = eo_dat_srv.get_value();
+
+ data<sch::Array<sch::Int16>,encode::NativeRaw> val{4};
+ val.at(0u).set(5);
+ val.at(1u).set(3);
+ val.at(2u).set(-6);
+ val.at(3u).set(1);
+ id<sch::Array<sch::Int16>> id_val{0u};
+ auto eo_send = dat_srv->send(val, id_val);
+ if(eo_send.is_error()){
+ return std::move(eo_send.get_error());
+ }
- print_value<<<dim3(2),dim3(2),0,hipStreamDefault>>>(*(v.get_device_data()));
+ auto eo_dfind = dat_srv->find(id_val);
+ if(eo_dfind.is_error()){
+ return std::move(eo_dfind.get_error());
+ }
+ auto dfind = eo_dfind.get_value();
+
+ auto& v = dfind();
+
+ print_array_vals<<<dim3(2),dim3(2),0,hipStreamDefault>>>(*(v.get_device_data()));
+ }
return make_void();
}