Skip to content

Commit

Permalink
Merge branch 'dev'
Browse files Browse the repository at this point in the history
  • Loading branch information
risi-kondor committed Oct 10, 2023
2 parents 9a5d016 + 6e1aef1 commit 6380b8c
Show file tree
Hide file tree
Showing 37 changed files with 3,264 additions and 31 deletions.
21 changes: 18 additions & 3 deletions common.txt
Original file line number Diff line number Diff line change
@@ -1,6 +1,14 @@
include $(ROOTDIR)/config.txt

CNINE_INCLUDES=-I$(CNINE_ROOT)/include
CNINE_INCLUDES=-I$(CNINE_ROOT)
CNINE_INCLUDES+=-I$(CNINE_ROOT)/include
CNINE_INCLUDES+=-I$(CNINE_ROOT)/modules
CNINE_INCLUDES+=-I$(CNINE_ROOT)/algorithms
CNINE_INCLUDES+=-I$(CNINE_ROOT)/combinatorial
CNINE_INCLUDES+=-I$(CNINE_ROOT)/containers
CNINE_INCLUDES+=-I$(CNINE_ROOT)/math
CNINE_INCLUDES+=-I$(CNINE_ROOT)/utility
CNINE_INCLUDES+=-I$(CNINE_ROOT)/wrappers
CNINE_INCLUDES+=-I$(CNINE_ROOT)/objects/scalar
CNINE_INCLUDES+=-I$(CNINE_ROOT)/objects/tensor
CNINE_INCLUDES+=-I$(CNINE_ROOT)/objects/ntensor
Expand All @@ -10,9 +18,12 @@ CNINE_INCLUDES+=-I$(CNINE_ROOT)/objects/tensor_array
CNINE_INCLUDES+=-I$(CNINE_ROOT)/objects/backendA -I$(CNINE_ROOT)/objects/backendB
CNINE_INCLUDES+=-I$(CNINE_ROOT)/objects/tensor_array/cell_maps -I$(CNINE_ROOT)/objects/tensor_array/cell_ops
CNINE_INCLUDES+=-I$(CNINE_ROOT)/objects/labeled
CNINE_INCLUDES+=-I$(CNINE_ROOT)/objects/labeled2

CNINE_CUDADIR=$(CNINE_ROOT)/cuda

SNOB2_INCLUDES=-I$(SNOB2_ROOT)/include -I$(SNOB2_ROOT)/combinatorial -I$(SNOB2_ROOT)/Sn

GELIB_INCLUDEDIR=$(ROOTDIR)/include
SO2DIR=$(ROOTDIR)/objects/SO2
SO3DIR=$(ROOTDIR)/objects/SO3
Expand All @@ -27,12 +38,12 @@ CFLAGS= -std=c++17 -O3 #-ferror-limit=1
INCLUDE= -I $(ROOTDIR)/include
LIBS= -lstdc++ -lm -lpthread

CFLAGS+=-DCNINE_RANGE_CHECKING

MACROS+=-DGELIB_COPY_WARNINGS
MACROS+=-DGELIB_MOVE_WARNINGS
MACROS+=-DGELIB_CONVERT_WARNINGS

MACROS=

ifdef GELIB_SO3CG_DEBUG
MACROS+=-D_GELIB_SO3CG_DEBUG
endif
Expand All @@ -45,3 +56,7 @@ ifdef WITH_CUBLAS
MACROS+=-D_WITH_CUBLAS
endif

ifdef EIGENDIR
CFLAGS+=-D_WITH_EIGEN
endif

8 changes: 8 additions & 0 deletions config.txt
Original file line number Diff line number Diff line change
Expand Up @@ -4,6 +4,14 @@ CNINE_ROOT=$(ROOTDIR)/../cnine/

#GELIB_SO3CG_DEBUG=t

#MACROS+=-DCNINE_COPY_WARNINGS
#MACROS+=-DCNINE_MOVE_WARNINGS
#MACROS+=-DCNINE_ASSIGN_WARNINGS
#MACROS+=-DCNINE_MOVEASSIGN_WARNINGS


SNOB2_ROOT=$(ROOTDIR)/../Snob2/


# ---- CUDA OPTIONS ----------------------------------

Expand Down
12 changes: 6 additions & 6 deletions include/GElibTimer.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -78,15 +78,15 @@ namespace GElib{

~CGproductTimer(){
auto elapsed=chrono::duration<double,std::milli>(chrono::system_clock::now()-t0).count();
int Mflops=0;
if(elapsed>0) Mflops=(((float)n_ops)/elapsed/1000.0);

if(gelib_log){
(*gelib_log)("CGproduct("+to_string(l1)+","+to_string(l2)+","+to_string(l)+")[b="+
to_string(b)+",n1="+to_string(n1)+",n2="+to_string(n2)+",dev="+to_string(dev)+"] "+
to_string(elapsed)+" ms"+" ["+to_string((int)(((float)n_ops)/elapsed/1000.0))+" Mflops]");
if(elapsed>0)
gelib_log->ofs2<<l1<<","<<l2<<","<<l<<","<<b<<","<<n1<<","<<n2<<","<<dev<<","<<elapsed<<","<<
(int)(((float)n_ops)/elapsed/1000.0)<<endl;
else
gelib_log->ofs2<<l1<<","<<l2<<","<<l<<","<<b<<","<<n1<<","<<n2<<","<<dev<<","<<elapsed<<","<<0<<endl;
to_string(elapsed)+" ms"+" ["+to_string(Mflops)+" Mflops]");

gelib_log->ofs2<<l1<<","<<l2<<","<<l<<","<<b<<","<<n1<<","<<n2<<","<<dev<<","<<elapsed<<","<<Mflops<<endl;

}
}
Expand Down
File renamed without changes.
10 changes: 10 additions & 0 deletions objects/SO3/SO3part3_view.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -13,6 +13,7 @@
#define _SO3part3_view

//#include "CtensorB.hpp"
#include "GstridesB.hpp"
#include "Ctensor3_view.hpp"
#include "SO3part2_view.hpp"
//#include "SO3_CGbank.hpp"
Expand All @@ -33,6 +34,7 @@ namespace GElib{
typedef cnine::fill_pattern fill_pattern;
typedef cnine::Gdims Gdims;
typedef cnine::Gstrides Gstrides;
typedef cnine::GstridesB GstridesB;

//float* arr;
//float* arrc;
Expand Down Expand Up @@ -76,6 +78,14 @@ namespace GElib{
ac=arrc+l*s1;
}

SO3part3_view(float* _arr, const Gdims& _dims, const GstridesB& _strides, const int _coffs=1, const int _dev=0):
Ctensor3_view(_arr,_dims,_strides,_coffs,_dev){
assert(n1%2==1);
l=(n1-1)/2;
ar=arr+l*s1;
ac=arrc+l*s1;
}

SO3part3_view(const Ctensor3_view& x):
Ctensor3_view(x){
assert(n1%2==1);
Expand Down
29 changes: 18 additions & 11 deletions objects/SO3/functions/SO3part_addCGproductFn.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -57,7 +57,7 @@ namespace GElib{
assert(l>=abs(l1-l2) && l<=l1+l2);

int count=0; for(int i=-l1; i<=l1; i++) count+=std::min(l2,l-i)-std::max(-l2,-l-i)+(i<=l);
CGproductTimer(l1,l2,l,B,N1,N2,dev,B*count*N1*N2);
CGproductTimer timer(l1,l2,l,B,N1,N2,dev,B*count*N1*N2);

if(dev==0 && cnine::dev_selector.dev>0){
int sdev=cnine::dev_selector.dev;
Expand All @@ -66,10 +66,11 @@ namespace GElib{
gelib_log->error(__PRETTY_FUNCTION__,"Arguments of streaming operation must have regular strides. Skipping this operation."); return;}

int nb=std::min((cnine::dev_selector.max_mem<<18)/(2*_x.n1*_x.n2+2*_y.n1*_y.n2+2*_r.n1*_x.n2*_y.n2),_r.n0);
cout<<"nb="<<nb;
int rsize_per_batch=_r.n1*_x.n2*_y.n2;
//cout<<"nb="<<nb<<endl;
cnine::ArrayOnDevice<float> xbuf(2*nb*_x.n1*_x.n2);
cnine::ArrayOnDevice<float> ybuf(2*nb*_y.n1*_y.n2);
cnine::ArrayOnDevice<float> rbuf(2*nb*_r.n1*_x.n2*_y.n2);
cnine::ArrayOnDevice<float> rbuf(2*nb*rsize_per_batch);
cnine::Ctensor3_view xv(xbuf,nb,_x.n1,_x.n2,_x.s0,_x.s1,_x.s2,1,sdev);
cnine::Ctensor3_view yv(ybuf,nb,_y.n1,_y.n2,_y.s0,_y.s1,_y.s2,1,sdev);
cnine::Ctensor3_view rv(rbuf,nb,_r.n1,_x.n2*_y.n2,_r.s0,_r.s1,_r.s2,1,sdev);
Expand All @@ -78,24 +79,30 @@ namespace GElib{
cu_stream stream;
//cudaStream_t stream;
//CUDA_SAFE(cudaStreamCreate(&stream));
cout<<"Streaming packets: "<<cnine::roundup(_r.n0,nb)/nb<<" "<<nb<<endl;
for(int i=0; i<cnine::roundup(_r.n0,nb)/nb; i++){

int _nb=std::max(nb,_x.n0-i*nb);
cout<<i<<" "<<i*nb*_r.s0+2*_offs<<endl;
int _nb=std::min(nb,_x.n0-i*nb);
if(nb==0) continue;
if(_nb<nb){
xv.n0=_nb;
yv.n0=_nb;
rv.n0=_nb;
}

CUDA_SAFE(cudaMemcpyAsync(xbuf, _x.arr+i*nb*_x.s0, _nb*_x.s0*sizeof(float), cudaMemcpyHostToDevice, stream));
CUDA_SAFE(cudaMemcpyAsync(ybuf, _y.arr+i*nb*_y.s0, _nb*_y.s0*sizeof(float), cudaMemcpyHostToDevice, stream));
CUDA_SAFE(cudaMemcpyAsync(rbuf, _r.arr+i*nb*_r.s0+2*_offs, _nb*_r.s0*sizeof(float), cudaMemcpyHostToDevice, stream));
SO3partB_addCGproduct_cu(rbuf,xbuf,ybuf,0,stream);
CUDA_SAFE(cudaMemcpyAsync(_r.arr+i*nb*_r.s0+2*_offs, rbuf, _nb*_r.s0*sizeof(float), cudaMemcpyDeviceToHost, stream));
//CUDA_SAFE(cudaMemcpyAsync(rbuf, _r.arr+i*nb*_r.s0+2*_offs, 2*_nb*rsize_per_batch*sizeof(float),
// cudaMemcpyHostToDevice, stream));
SO3partB_addCGproduct_cu(rv,xv,yv,0,stream);
//CUDA_SAFE(cudaMemcpyAsync(_r.arr+i*nb*_r.s0+2*_offs, rbuf, 2*_nb*rsize_per_batch*sizeof(float),
// cudaMemcpyDeviceToHost, stream));

}
//CUDA_SAFE(cudaStreamSynchronize(stream));
//CUDA_SAFE(cudaStreamDestroy(stream));
#endif

return;
}

if(dev==0){
Expand All @@ -120,8 +127,8 @@ namespace GElib{
});
}

else CUDA_STREAM(SO3partB_addCGproduct_cu(_r,_x,_y,_offs,stream));

if(dev==1) CUDA_STREAM(SO3partB_addCGproduct_cu(_r,_x,_y,_offs,stream));
}


Expand Down
2 changes: 1 addition & 1 deletion objects/SO3n/SO3bipartArrayView.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -96,7 +96,7 @@ namespace GElib{

operator cnine::Ctensor4_view() const{
return cnine::Ctensor4_view(arr.template ptr_as<RTYPE>(),{getb()*getN(),dims(-3),dims(-2),dims(-1)},
{2*strides(-4),2*strides(-3),2*strides(-2),2*strides(-1)},1,device());
cnine::GstridesB(2*strides(-4),2*strides(-3),2*strides(-2),2*strides(-1)),1,device());
}


Expand Down
2 changes: 1 addition & 1 deletion objects/SO3n/SO3bipartView.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -67,7 +67,7 @@ namespace GElib{

operator cnine::Ctensor4_view() const{
return cnine::Ctensor4_view(arr.template ptr_as<RTYPE>(),{dims[0],dims[1],dims[2],dims[3]},
{2*strides[0],2*strides[1],2*strides[2],2*strides[3]},1,device());
cnine::GstridesB(2*strides[0],2*strides[1],2*strides[2],2*strides[3]),1,device());
}


Expand Down
4 changes: 2 additions & 2 deletions objects/SO3n/SO3partArrayView.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -92,12 +92,12 @@ namespace GElib{

operator cnine::Ctensor3_view() const{
return cnine::Ctensor3_view(arr.template ptr_as<RTYPE>(),{getb()*getN(),dims(-2),dims(-1)},
{2*strides(-3),2*strides(-2),2*strides(-1)},1,device());
cnine::GstridesB(2*strides(-3),2*strides(-2),2*strides(-1)),1,device());
}

operator SO3part3_view() const{
return SO3part3_view(arr.template ptr_as<RTYPE>(),{getb()*getN(),dims(-2),dims(-1)},
{2*strides(-3),2*strides(-2),2*strides(-1)},1,device());
cnine::GstridesB(2*strides(-3),2*strides(-2),2*strides(-1)),1,device());
}


Expand Down
4 changes: 2 additions & 2 deletions objects/SO3n/SO3partView.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -62,12 +62,12 @@ namespace GElib{

operator cnine::Ctensor3_view() const{
return cnine::Ctensor3_view(arr.template ptr_as<RTYPE>(),{dims[0],dims[1],dims[2]},
{2*strides[0],2*strides[1],2*strides[2]},1,device());
cnine::GstridesB(2*strides[0],2*strides[1],2*strides[2]),1,device());
}

operator SO3part3_view() const{
return SO3part3_view(arr.template ptr_as<RTYPE>(),{dims[0],dims[1],dims[2]},
{2*strides[0],2*strides[1],2*strides[2]},1,device());
cnine::GstridesB(2*strides[0],2*strides[1],2*strides[2]),1,device());
}


Expand Down
12 changes: 7 additions & 5 deletions objects/SO3n/tests/testSO3partArray_Streaming.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -14,18 +14,20 @@ int main(int argc, char** argv){
GElibSession session;
cout<<endl;

int b=2;
int l=2;
int nc=2;
Gdims adims({2});
int b=200;
int l=4;
int nc=32;
Gdims adims({40,40});

SO3partArray<float> u=SO3partArray<float>::gaussian(b,adims,l,nc);
SO3partArray<float> v=SO3partArray<float>::gaussian(b,adims,l,nc);
//printl("u",u)<<endl;
//printl("v",v)<<endl;

SO3partArray<float> w=StreamingCGproduct(u,v,2);
//xsSO3partArray<float> w=CGproduct(u,v,2);
//cout<<w<<endl;
SO3partArray<float> ws=StreamingCGproduct(u,v,2);
//cout<<ws<<endl;

//cout<<DiagCGproduct(u,v,2)<<endl;

Expand Down
125 changes: 125 additions & 0 deletions objects/prod/CGprodBasis.hpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,125 @@
// This file is part of GElib, a C++/CUDA library for group
// equivariant tensor operations.
//
// Copyright (c) 2023, Imre Risi Kondor
//
// This Source Code Form is subject to the terms of the Mozilla
// Public License v. 2.0. If a copy of the MPL was not distributed
// with this file, You can obtain one at http://mozilla.org/MPL/2.0/.

#ifndef _CGprodBasis
#define _CGprodBasis

#include "CGprodBasisObj.hpp"


namespace GElib{


template<typename GROUP>
class CGprodBasis{
public:

typedef typename GROUP::IrrepIx _IrrepIx;
typedef CGprodBasisObj<GROUP> OBJ;


OBJ* obj;

CGprodBasis(OBJ* _obj):
obj(_obj){}

CGprodBasis(OBJ& _obj):
obj(&_obj){}

CGprodBasis(const _IrrepIx& ix):
CGprodBasis(GROUP::space(ix)){}



public: // ---- Access ------------------------------------------------------------------------------------


bool is_leaf() const{
return obj->is_leaf();
}

bool is_isomorphic(const CGprodBasis& y) const{
return obj->is_isomorphic(*y.obj);
}

CGprodBasis left() const{
GELIB_ASSRT(!is_leaf());
return obj->left;
}

CGprodBasis right() const{
GELIB_ASSRT(!is_leaf());
return obj->right;
}

Gtype<GROUP> get_tau() const{
return obj->get_tau();
}


public: // ---- Transformations to other bases ------------------------------------------------------------


CGprodBasis shift_left() const{
return obj->shift_left();
}

CGprodBasis shift_right() const{
return obj->shift_right();
}

CGprodBasis standard_form() const{
return obj->standard_form();
}

CGprodBasis reverse_standard_form() const{
return obj->reverse_standard_form();
}

const EndMap<GROUP,double>& standardizing_map() const{
return obj->standardizing_map();
}

const EndMap<GROUP,double>& swap_map() const{
return obj->swap_map();
}

const EndMap<GROUP,double>& transpose_last_map() const{
return obj->transpose_last_map();
}



public: // ---- I/O ---------------------------------------------------------------------------------------


string repr() const{
return obj->repr();
}

string str(const string indent="") const{
return obj->str(indent);
}

friend ostream& operator<<(ostream& stream, const CGprodBasis& x){
stream<<x.str(); return stream;
}


};


template<typename GROUP>
inline CGprodBasis<GROUP> operator*(const CGprodBasis<GROUP>& x, const CGprodBasis<GROUP>& y){
return CGprodBasis(GROUP::space(x.obj,y.obj));
}

}

#endif
Loading

0 comments on commit 6380b8c

Please sign in to comment.