Putting this into a repo where Christian and I can work on it

(others can maybe join later)

KokkosProxies.hpp

@@ -0,0 +1,102 @@
+/*
+ * KokkosProxies.hpp
+ *
+ *  Created on: Sep 9, 2019
+ *      Author: bjoo
+ */
+
+#ifndef KOKKOSPROXIES_HPP_
+#define KOKKOSPROXIES_HPP_
+
+namespace Foo {
+namespace Bar {
+
+// Proxy Range Policy
+struct RangePolicy {
+	size_t start_idx;
+	size_t num_items;
+
+	RangePolicy( size_t start_idx_, size_t num_items_) : start_idx(start_idx_), num_items(num_items_) {}
+	inline
+	size_t begin() { return start_idx; }
+
+	// Past last item.
+	inline
+	size_t end() { return num_items; }
+};
+
+
+// Proxy parallel for closure
+template< class FunctorType>
+class ParallelFor {
+public:
+	using Policy = Foo::Bar::RangePolicy;
+
+	FunctorType m_functor;
+	Policy m_policy;
+
+	inline
+	void execute() const {
+		launch(*this);
+	}
+
+
+	ParallelFor( const FunctorType & arg_functor ,
+			const Policy   &     arg_policy )
+	: m_functor( arg_functor )
+	, m_policy(  arg_policy )
+	{ }
+
+
+
+};
+
+// Proxy Launcher
+template<class Driver>
+void launch(Driver driver_in) {
+	std::cerr << "In sycl_launch_copy" << std::endl;
+
+	// cl::sycl::queue* q = driver_in.m_policy.space().impl_internal_space_instance()->m_queue;
+	cl::sycl::queue* q = Kokkos::Experimental::SYCL().impl_internal_space_instance()->m_queue;
+	std::cerr << "Queue pointer is: " << (unsigned long) q << std::endl;
+
+	std::cerr << "range=" << driver_in.m_policy.end()-driver_in.m_policy.begin() << std::endl;
+	std::cerr << "driver_in.ptr_d = " << (unsigned long)(driver_in.m_functor.ptr_d) << std::endl;
+
+	q->submit([&](cl::sycl::handler& cgh) {
+		cl::sycl::stream out(1024,256,cgh);
+		cgh.parallel_for (
+				cl::sycl::range<1>(driver_in.m_policy.end()-driver_in.m_policy.begin()),
+				[=] (cl::sycl::id<1> item) {
+			size_t idx = item[0];
+			if (idx == 2 ) { // stop threads overwriting
+				out << "idx = " << idx << " PF ptr_d = " << (unsigned long)driver_in.m_functor.ptr_d << cl::sycl::endl;
+			}
+			driver_in.m_functor(idx, out);
+		});
+	});
+	q->wait_and_throw();
+}
+
+
+// Proxy ParallelFor Dispatch call
+template <class FunctorType>
+inline void my_parallel_for_2(const size_t work_count, const FunctorType& functor,
+		const std::string& str = "") {
+
+	//Foo::Bar::RangePolicy policy(0,work_count);
+	Foo::Bar::ParallelFor<FunctorType> closure(functor,
+			Foo::Bar::RangePolicy(0, work_count));
+
+
+	std::cerr << "Calling execute" <<std::endl;
+
+	closure.execute();
+}
+
+} // namespace Bar
+} // Namespace Foo
+
+
+
+#endif /* KOKKOSPROXIES_HPP_ */

Makefile

@@ -0,0 +1,47 @@
+KOKKOS_DEVICES=SyCL
+KOKKOS_CUDA_OPTIONS=enable_lambda
+KOKKOS_ARCH = "None"
+
+
+MAKEFILE_PATH := $(subst Makefile,,$(abspath $(lastword $(MAKEFILE_LIST))))
+KOKKOS_PATH=/home/bjoo/KokkosSyCL/kokkos
+ifndef KOKKOS_PATH
+  KOKKOS_PATH = $(MAKEFILE_PATH)../..
+endif
+
+KOKKOS_OPTIONS=disable_deprecated_code
+
+SRC = $(wildcard $(MAKEFILE_PATH)*.cpp)
+HEADERS = $(wildcard $(MAKEFILE_PATH)*.hpp)
+
+vpath %.cpp $(sort $(dir $(SRC)))
+
+default: build
+	echo "Start Build"
+
+EXE = bytes_and_flops.host
+CXX=clang++
+CXXFLAGS ?= -O3 -g
+override CXXFLAGS += -I$(MAKEFILE_PATH)
+
+DEPFLAGS = -M
+LINK = ${CXX}
+LINKFLAGS =
+
+OBJ = $(notdir $(SRC:.cpp=.o))
+LIB =
+
+include $(KOKKOS_PATH)/Makefile.kokkos
+
+build: $(EXE)
+
+$(EXE): $(OBJ) $(KOKKOS_LINK_DEPENDS)
+	$(LINK) $(KOKKOS_LDFLAGS) $(LINKFLAGS) $(EXTRA_PATH) $(OBJ) $(KOKKOS_LIBS) $(LIB) -o $(EXE)
+
+clean: kokkos-clean
+	rm -f *.o *.cuda *.host
+
+# Compilation rules
+
+%.o:%.cpp $(KOKKOS_CPP_DEPENDS) $(HEADERS)
+	$(CXX) $(KOKKOS_CPPFLAGS) $(KOKKOS_CXXFLAGS) $(CXXFLAGS) $(EXTRA_INC) -c $< -o $(notdir $@)

main.cpp

@@ -0,0 +1,77 @@
+#include <Kokkos_Core.hpp>
+struct functor {
+  functor( void * ptr_d_ ) : ptr_d((unsigned long)ptr_d_), mult(2.0) {}
+
+  double mult;   
+  unsigned long ptr_d=0xdeadbeef;
+
+
+  inline
+  void operator()(const int i,  cl::sycl::stream out) const {
+    double *dptr = (double *)ptr_d;
+
+    if ( dptr == 0x0 ) {
+      if ( i  == 2 ) { 
+	out << "Id = " << i << " BARF!!!" << cl::sycl::endl;
+      }
+    }
+    else {
+      dptr[i] = mult*(double)i+1.5;
+    }
+  }
+};
+
+void copyAndPrint(void *ptr_d, cl::sycl::queue* q, int N);
+
+#include "KokkosProxies.hpp"
+
+int main(int argc, char *argv[])
+{
+	Kokkos::initialize(argc,argv);
+
+	cl::sycl::queue* q = Kokkos::Experimental::SYCL().impl_internal_space_instance()->m_queue;
+	auto context = q->get_context();
+	auto  device = q->get_device();
+
+	const int N = 15;
+
+	void* ptr_d=cl::sycl::malloc_device(N*sizeof(double),device,context);
+	std::cout << " q ptr is : " << (unsigned long)q << std::endl;
+	std::cout << " ptr_d is : " << (unsigned long)ptr_d << std::endl;
+	std::cout << " Calling q.submit() " << std::endl;
+
+	functor f(ptr_d);
+	f.mult=4.0;
+	Foo::Bar::my_parallel_for_2(N,f);
+	copyAndPrint(ptr_d,q,N);
+
+	std::cout << "Kokkos::parallel_for " << std::endl;
+	f.mult = 6.0;
+	Kokkos::parallel_for(N,f);
+	Kokkos::fence();
+	copyAndPrint(ptr_d,q,N);
+
+	cl::sycl::free((void *)ptr_d,context);
+	Kokkos::finalize();
+}
+
+void copyAndPrint(void *ptr_d, cl::sycl::queue* q, int N) {
+
+
+	cl::sycl::buffer<double,1> hbuf(N);
+	q->submit([&](cl::sycl::handler& cgh) {
+	     auto access = hbuf.get_access<cl::sycl::access::mode::write>(cgh);
+	     cgh.single_task<class copy>([=]{
+	       for(int i=0; i<N; i++) {
+                 access[i] = ((double *)ptr_d)[i];
+	       }
+	     });
+	 });
+
+	auto haccess = hbuf.get_access<cl::sycl::access::mode::read>();
+	for(int i=0; i < N; ++i) {
+	  printf("%d %lf\n", i, haccess[i]);
+	}
+
+}
+