portelli/Grid
1
0
Fork 0
mirror of https://github.com/paboyle/Grid.git synced 2024-11-09 23:45:36 +00:00
Code Releases Activity

Making SYCL happy

Browse Source
This commit is contained in:
Peter Boyle 2023-09-26 13:19:42 -07:00
parent b8a7004365
commit 3d437c5cc4
2 changed files with 6 additions and 4 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
Grid/communicator/SharedMemoryMPI.cc
View File

@ -604,8 +604,8 @@ void GlobalSharedMemory::SharedMemoryAllocate(uint64_t bytes, int flags)
#ifdef GRID_SYCL_LEVEL_ZERO_IPC #ifdef GRID_SYCL_LEVEL_ZERO_IPC
typedef struct { int fd; pid_t pid ; ze_ipc_mem_handle_t ze; } clone_mem_t; typedef struct { int fd; pid_t pid ; ze_ipc_mem_handle_t ze; } clone_mem_t;
auto zeDevice = cl::sycl::get_native<cl::sycl::backend::level_zero>(theGridAccelerator->get_device()); auto zeDevice = cl::sycl::get_native<cl::sycl::backend::ext_oneapi_level_zero>(theGridAccelerator->get_device());
auto zeContext = cl::sycl::get_native<cl::sycl::backend::level_zero>(theGridAccelerator->get_context()); auto zeContext = cl::sycl::get_native<cl::sycl::backend::ext_oneapi_level_zero>(theGridAccelerator->get_context());
ze_ipc_mem_handle_t ihandle; ze_ipc_mem_handle_t ihandle;
clone_mem_t handle; clone_mem_t handle;

6
Grid/tensors/Tensor_Ta.h
View File

@ -90,10 +90,12 @@ template<class vtype,int N> accelerator_inline iVector<vtype,N> ProjectOnGroup(c
template<class vtype,int N, typename std::enable_if< GridTypeMapper<vtype>::TensorLevel == 0 >::type * =nullptr> template<class vtype,int N, typename std::enable_if< GridTypeMapper<vtype>::TensorLevel == 0 >::type * =nullptr>
accelerator_inline iMatrix<vtype,N> ProjectOnGroup(const iMatrix<vtype,N> &arg) accelerator_inline iMatrix<vtype,N> ProjectOnGroup(const iMatrix<vtype,N> &arg)
{ {
typedef typename iMatrix<vtype,N>::scalar_type scalar;
// need a check for the group type? // need a check for the group type?
iMatrix<vtype,N> ret(arg); iMatrix<vtype,N> ret(arg);
vtype nrm; vtype nrm;
vtype inner; vtype inner;
scalar one(1.0);
for(int c1=0;c1<N;c1++){ for(int c1=0;c1<N;c1++){
// Normalises row c1 // Normalises row c1
@ -102,7 +104,7 @@ accelerator_inline iMatrix<vtype,N> ProjectOnGroup(const iMatrix<vtype,N> &arg)
inner += innerProduct(ret._internal[c1][c2],ret._internal[c1][c2]); inner += innerProduct(ret._internal[c1][c2],ret._internal[c1][c2]);
nrm = sqrt(inner); nrm = sqrt(inner);
nrm = 1.0/nrm; nrm = one/nrm;
for(int c2=0;c2<N;c2++) for(int c2=0;c2<N;c2++)
ret._internal[c1][c2]*= nrm; ret._internal[c1][c2]*= nrm;
@ -127,7 +129,7 @@ accelerator_inline iMatrix<vtype,N> ProjectOnGroup(const iMatrix<vtype,N> &arg)
inner += innerProduct(ret._internal[c1][c2],ret._internal[c1][c2]); inner += innerProduct(ret._internal[c1][c2],ret._internal[c1][c2]);
nrm = sqrt(inner); nrm = sqrt(inner);
nrm = 1.0/nrm; nrm = one/nrm;
for(int c2=0;c2<N;c2++) for(int c2=0;c2<N;c2++)
ret._internal[c1][c2]*= nrm; ret._internal[c1][c2]*= nrm;
} }