Commit 0e143080 authored by Pawel Markiewicz's avatar Pawel Markiewicz

FIRST RELEASE

parents
Copyright (c) 2016, University College London, United Kingdom
All rights reserved.
Redistribution and use in source and binary forms, with or without
modification, are permitted provided that the following conditions are met:
* Redistributions of source code must retain the above copyright notice, this
list of conditions and the following disclaimer.
* Redistributions in binary form must reproduce the above copyright notice,
this list of conditions and the following disclaimer in the documentation
and/or other materials provided with the distribution.
* Neither the name of [project] nor the names of its
contributors may be used to endorse or promote products derived from
this software without specific prior written permission.
THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
\ No newline at end of file
graft nipet/auxdata
graft nipet/src
graft nipet/img/src
graft nipet/lm/src
include nipet/def.h
include nipet/Makefile
include nipet/img/Makefile
include nipet/lm/Makefile
OBJDIR=obj
SRCDIR=src
INCDIR=.
CC=nvcc
NFLAGS=-arch=sm_35 --ptxas-options=-v -Xcompiler '-m64','-fPIC','-DNDEBUG','-O3','-Wall','-Wstrict-prototypes'
CFLAGS=-I/usr/include/python2.7 -I$(INCDIR)
LDFLAGS=-L/usr/local/lib -L/usr/local/cuda/lib64 -lcudart
DEPS = $(INCDIR)/def.h $(SRCDIR)/aux.h $(SRCDIR)/norm.h $(SRCDIR)/auxmath.h
OBJ = $(OBJDIR)/aux_module.o $(OBJDIR)/auxmath.o $(OBJDIR)/aux.o $(OBJDIR)/norm.o
ODIR = ./
$(OBJDIR)/%.o: $(SRCDIR)/%.cu $(DEPS)
$(CC) $(NFLAGS) -c -o $@ $< $(CFLAGS)
$(ODIR)mmr_auxe.so: $(OBJ)
nvcc -shared -o $@ $^ $(LDFLAGS)
$(OBJ): | $(OBJDIR)
$(OBJDIR):
mkdir -p ./$(OBJDIR)
\ No newline at end of file
# init the package folder
import mmraux
import mmrnorm
import lm
import img
import prj
import sct
\ No newline at end of file
#include <stdio.h>
#ifndef _DEF_H_
#define _DEF_H_
#define EX_PRINT_INFO 0
#define RD2MEM 0
// device
#define BTHREADS 10
#define NTHREADS 256
#define TOTHRDS (BTHREADS*NTHREADS)
#define ITIME 1000 //integration time
#define BTPTIME 100 //time period for bootstrapping
#define MVTIME 1000
#define VTIME 2 // 2**VTIME = time resolution for PRJ VIEW [s]
#define MXNITAG 7200 //max number of time tags <nitag> to avoid out of memory errors
#define TOT_BINS_S1 354033792 //344*252*4084
//344*252*837
#define TOT_BINS 72557856
#define NSTREAMS 32 // # CUDA streams
#define ELECHNK (402653184/NSTREAMS) //2^28 = 268435456 int elements to make up 1GB 134217728 (402653184 = 2^28+2^27 => 1.5G), 536870912
//=== LM bit fields/masks ===
// mask for time bits
#define mMR_TMSK (0x1fffffff)
// check if time tag
#define mMR_TTAG(w) ( (w>>29) == 4 )
//for randoms
#define mxRD 60 //maximum ring difference
#define CFOR 20 //number of iterations for crystals transaxially
#define SPAN 11
#define NRINGS 64
#define nCRS 504
#define nCRSR 448 // number of active crystals
#define NSBINS 344
#define NSANGLES 252
#define NSBINANG 86688 //NSBINS*NSANGLES
#define NSINOS 4084
#define NSINOS11 837
#define SEG0 127
#define NBUCKTS 224 //purposely too large (should be 224 = 28*8)
#define AW 68516 //number of active bins in 2D sino
#define NLI2R 2074
//coincidence time window in pico-seconds
#define CWND = 5859.38
typedef struct{
char*fname;
off_t *atag;
int *btag;
int *ele4chnk;
int *ele4thrd;
size_t ele;
int nchnk;
int nitag;
int toff;
int flgs; //write out sinos in span-11
int span; //choose span (1, 11 or SSRB)
int nfrm; //ouput dynamic sinos in span-11
int flgf; //do fansums calculatins and output by randoms estimation
int nfrm2;
short *t2dfrm;
int frmoff; //frame offset to account for the splitting of the dynamic data into two
int last_ttag;
int tstart;
int tstop;
int tmidd;
} LMprop; //properties of LM data file and its breaking up into chunks of data.
#define PI 3.1415926535f
//temporary...
#define L21 0.001f // threshold for special case when finding Siddon intersections
#define TA1 0.7885139f // angle threshold 1 for Siddon calculations ~ PI/4
#define TA2 -0.7822831f // angle threshold 2 for Siddon calculations ~-PI/4
#define N_TV 407 // max number of voxels intersections with a ray (t)
#define N_TT 10 // number of constants precaluclated and saved for proper axial calulations
#define UV_SHFT 9 // shift when representing 2 voxel indx in one float variable
//actual radius (radious + DOI)
#define R_RING (32.8 + 0.67)
#define IR_RNG 0.0298775
#define R_2 1120.2409f
#define TFOV2 890.00f //not sure why it is so...
//#define RS 35.00f //radius for scatter calculation, approx. SZ_VOXY*SE_IMX/2
#define SZ_IMZ 127
#define SZ_IMP 128 // axial image size padded
#define HSZ_IMZ 63.5
#define SZ_IMX 320
#define SZ_IMY 320
#define SE_IMX 344
#define SE_IMY 344
//<><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><>
#define SZ_RING 0.40625f
#define aXR 0.40625f
#define SZ_VOXY 0.208626f
#define _SZVXY 4.7932664f
#define SZ_VOXZ 0.203125f
#define ISZ_VOXZ 4.9230771f
//crystal angle
#define aLPHA ((2*PI)/nCRS)
#endif // end of _DEF_H_
// //============== SCATTER, KN LUT =====================
// //maximum scattering angle due to Elow
// //#define UPSMX 0.64744709f
// //279 //400 //((1-COSUPSMX)/COSSTP)+1
// // //cos step (instead of angle sample cosines)
// // #define COSSTP 1e3
// //cos(UPSMX)
// // #define COSUPSMX 0.74//0.722 //LLD = E511/(2-cos(upsmx)) //0.601 //
// //number of samples
// #define NCOS 256
// //number of samples per scattering patch length. it is the power of 2. 2**LSCT2 = patch length
// #define LSCT2 2
// //detector surface
// #define SRFCRS 0.1695112f //2*math.pi/NCRS*(32.8 + 0.67)*0.40625
// //#define SRFCRS_R2 0.00015373925f
// //====================================================
\ No newline at end of file
OBJDIR=obj
SRCDIR=src
CC=nvcc
NFLAGS=-arch=sm_35 --ptxas-options=-v -Xcompiler '-m64','-fPIC','-DNDEBUG','-O3','-Wall','-Wstrict-prototypes'
CFLAGS=-I/usr/include/python2.7 -I$(INCDIR) -I$(INCDIR2)
LDFLAGS=-L/usr/local/lib -L/usr/local/cuda/lib64 -lcudart
DEPS = $(SRCDIR)/rsmpl.h
OBJ = $(OBJDIR)/img_module.o $(OBJDIR)/rsmpl.o
ODIR = ./
$(OBJDIR)/%.o: $(SRCDIR)/%.cu $(DEPS)
$(CC) $(NFLAGS) -c -o $@ $< $(CFLAGS)
$(ODIR)improc.so: $(OBJ)
nvcc -shared -o $@ $^ $(LDFLAGS)
$(OBJ): | $(OBJDIR)
$(OBJDIR):
mkdir -p ./$(OBJDIR)
\ No newline at end of file
# init the package folder
import mmrimg
\ No newline at end of file
This diff is collapsed.
#include <Python.h>
#include <numpy/arrayobject.h>
#include <stdlib.h>
#include "rsmpl.h"
#include "auxmath.h"
//=== PYTHON STUFF ===
//--- Docstrings
static char module_docstring[] =
"This module provides GPU routines for (mostly PET) image processing.";
static char rsmpl_docstring[] =
"Does rigid body transformation with very fine sampling.";
//---
//--- Available functions
static PyObject *img_resample(PyObject *self, PyObject *args);
/* Module specification */
static PyMethodDef module_methods[] = {
{"resample", img_resample, METH_VARARGS, rsmpl_docstring},
{NULL, NULL, 0, NULL}
};
//---
//--- Initialize the module
PyMODINIT_FUNC initimproc(void) //it HAS to be init______ and then the name of the shared lib.
{
PyObject *m = Py_InitModule3("improc", module_methods, module_docstring);
if (m == NULL)
return;
/* Load NumPy functionality. */
import_array();
}
//=======================
//======================================================================================
// P R O C E S I N G I M A G E D A T A
//--------------------------------------------------------------------------------------
static PyObject *img_resample(PyObject *self, PyObject *args)
{
// transformation matrix
PyObject * o_A;
// Structure for constants
Cimg Cim;
//Dictionary of image constants
PyObject * o_Cim;
//original image (to be transformed)
PyObject * o_imo;
/* Parse the input tuple */
if ( !PyArg_ParseTuple(args, "OOO", &o_imo, &o_A, &o_Cim))
return NULL;
//the dictionary of constants
PyObject* pd_vxsox = PyDict_GetItemString(o_Cim, "VXSOx");
Cim.VXSOx = (float) PyFloat_AsDouble(pd_vxsox);
PyObject* pd_vxsoy = PyDict_GetItemString(o_Cim, "VXSOy");
Cim.VXSOy = (float) PyFloat_AsDouble(pd_vxsoy);
PyObject* pd_vxsoz = PyDict_GetItemString(o_Cim, "VXSOz");
Cim.VXSOz = (float) PyFloat_AsDouble(pd_vxsoz);
PyObject* pd_vxnox = PyDict_GetItemString(o_Cim, "VXNOx");
Cim.VXNOx = (short) PyInt_AS_LONG(pd_vxnox);
PyObject* pd_vxnoy = PyDict_GetItemString(o_Cim, "VXNOy");
Cim.VXNOy = (short) PyInt_AS_LONG(pd_vxnoy);
PyObject* pd_vxnoz = PyDict_GetItemString(o_Cim, "VXNOz");
Cim.VXNOz = (short) PyInt_AS_LONG(pd_vxnoz);
PyObject* pd_offox = PyDict_GetItemString(o_Cim, "OFFOx");
Cim.OFFOx = (float) PyFloat_AsDouble(pd_offox);
PyObject* pd_offoy = PyDict_GetItemString(o_Cim, "OFFOy");
Cim.OFFOy = (float) PyFloat_AsDouble(pd_offoy);
PyObject* pd_offoz = PyDict_GetItemString(o_Cim, "OFFOz");
Cim.OFFOz = (float) PyFloat_AsDouble(pd_offoz);
PyObject* pd_vxsrx = PyDict_GetItemString(o_Cim, "VXSRx");
Cim.VXSRx = (float) PyFloat_AsDouble(pd_vxsrx);
PyObject* pd_vxsry = PyDict_GetItemString(o_Cim, "VXSRy");
Cim.VXSRy = (float) PyFloat_AsDouble(pd_vxsry);
PyObject* pd_vxsrz = PyDict_GetItemString(o_Cim, "VXSRz");
Cim.VXSRz = (float) PyFloat_AsDouble(pd_vxsrz);
PyObject* pd_vxnrx = PyDict_GetItemString(o_Cim, "VXNRx");
Cim.VXNRx = (short) PyInt_AS_LONG(pd_vxnrx);
PyObject* pd_vxnry = PyDict_GetItemString(o_Cim, "VXNRy");
Cim.VXNRy = (short) PyInt_AS_LONG(pd_vxnry);
PyObject* pd_vxnrz = PyDict_GetItemString(o_Cim, "VXNRz");
Cim.VXNRz = (short) PyInt_AS_LONG(pd_vxnrz);
PyObject* pd_offrx = PyDict_GetItemString(o_Cim, "OFFRx");
Cim.OFFRx = (float) PyFloat_AsDouble(pd_offrx);
PyObject* pd_offry = PyDict_GetItemString(o_Cim, "OFFRy");
Cim.OFFRy = (float) PyFloat_AsDouble(pd_offry);
PyObject* pd_offrz = PyDict_GetItemString(o_Cim, "OFFRz");
Cim.OFFRz = (float) PyFloat_AsDouble(pd_offrz);
PyObject *p_A = PyArray_FROM_OTF(o_A, NPY_FLOAT32, NPY_IN_ARRAY);
PyObject *p_imo = PyArray_FROM_OTF(o_imo, NPY_FLOAT32, NPY_IN_ARRAY);
/* If that didn't work, throw an exception. */
if (p_A == NULL || p_imo == NULL) {
Py_XDECREF(p_A);
Py_XDECREF(p_imo);
return NULL;
}
float *A = (float*)PyArray_DATA(p_A);
float *imo = (float*) PyArray_DATA(p_imo);
// for (int i=0; i<12; i++) printf("A[%d] = %f\n",i,A[i] );
//=================================================================
float *imr = rsmpl(imo, A, Cim);
//=================================================================
printf("i> new image (x,y,z) = (%d,%d,%d)\n voxel size: (%6.4f, %6.4f, %6.4f)\n", Cim.VXNRx, Cim.VXNRy, Cim.VXNRz, Cim.VXSRx, Cim.VXSRy, Cim.VXSRz);
npy_intp dims[3];
dims[2] = Cim.VXNRx;
dims[1] = Cim.VXNRy;
dims[0] = Cim.VXNRz;
PyArrayObject *p_imr = (PyArrayObject *)PyArray_SimpleNewFromData(3, dims, NPY_FLOAT32, imr);
// //--- form output tuples
// PyObject *tuple_out = PyTuple_New(2);
// PyTuple_SetItem(tuple_out, 0, Py_BuildValue("i", 23));
// PyTuple_SetItem(tuple_out, 1, PyArray_Return(p_imr));
// //---
//Clean up:
Py_DECREF(p_A);
Py_DECREF(p_imo);
return PyArray_Return(p_imr); //tuple_out;
}
#include "rsmpl.h"
void HandleError( cudaError_t err, const char *file, int line ){
if (err != cudaSuccess) {
printf( "%s in %s at line %d\n", cudaGetErrorString( err ), file, line );
exit( EXIT_FAILURE );
}
}
//..................................................................................................................................
__constant__ float cA[12];
__global__
void d_rsmpl(float *imr,
const float *imo,
Cimg Cim){
// extern __shared__ float s[];
int ib = blockIdx.x + blockIdx.y*gridDim.x + blockIdx.z*gridDim.x*gridDim.y;
//int it = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*blockDim.x*blockDim.y;
float x = (blockIdx.x*Cim.VXSOx + Cim.OFFOx) + Cim.VXSOx/VDIV*(0.5 + threadIdx.x);
float y = -(blockIdx.y*Cim.VXSOy + Cim.OFFOy) - Cim.VXSOy/VDIV*(0.5 + threadIdx.y);
float z = blockIdx.z*Cim.VXSOz + Cim.OFFOz + Cim.VXSOz/VDIV*(0.5 + threadIdx.z);
float xp = cA[0]*x + cA[1]*y + cA[2]*z + cA[3];
float yp = cA[4]*x + cA[5]*y + cA[6]*z + cA[7];
float zp = cA[8]*x + cA[9]*y + cA[10]*z + cA[11];
short u = roundf(-Cim.OFFRx/Cim.VXSRx) + floorf((xp)/Cim.VXSRx);
short v = roundf(-Cim.OFFRy/Cim.VXSRy) - ceilf ((yp)/Cim.VXSRy);
short w = roundf(-Cim.OFFRz/Cim.VXSRz) + floorf((zp)/Cim.VXSRz);
if ((u<Cim.VXNRx)&&(v<Cim.VXNRy)&&(w<Cim.VXNRz)&&(u>=0)&&(v>=0)&&(w>=0))
atomicAdd(imr + u+v*Cim.VXNRx+w*Cim.VXNRx*Cim.VXNRy, imo[ib]/(VDIV*VDIV*VDIV));
}
//::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::
float *rsmpl( float *imo,
float *A,
Cimg Cim)
{
float *d_imr;
HANDLE_ERROR( cudaMalloc(&d_imr, Cim.VXNRx*Cim.VXNRy*Cim.VXNRz*sizeof(float)) );
HANDLE_ERROR( cudaMemset(d_imr, 0, Cim.VXNRx*Cim.VXNRy*Cim.VXNRz*sizeof(float)) );
float *d_imo;
HANDLE_ERROR( cudaMalloc(&d_imo, Cim.VXNOx*Cim.VXNOy*Cim.VXNOz*sizeof(float)) );
HANDLE_ERROR( cudaMemcpy(d_imo, imo, Cim.VXNOx*Cim.VXNOy*Cim.VXNOz*sizeof(float), cudaMemcpyHostToDevice) );
cudaMemcpyToSymbol(cA, A, 12*sizeof(float));
// double * d_A;
// HANDLE_ERROR( cudaMalloc(&d_A, 12*sizeof(double)) );
// HANDLE_ERROR( cudaMemcpy(d_A, A, 12*sizeof(double), cudaMemcpyHostToDevice) );
//<><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><>
printf("i> calculating transformation with %d samples per voxel...", VDIV);
dim3 grid(Cim.VXNOx, Cim.VXNOy, Cim.VXNOz);
dim3 block(VDIV, VDIV, VDIV);
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
d_rsmpl<<<grid, block>>>(d_imr, d_imo, Cim);
cudaError_t error = cudaGetLastError();
if(error != cudaSuccess){printf("CUDA kernel for image resampling: error: %s\n", cudaGetErrorString(error)); exit(-1);}
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop);
cudaEventDestroy(start);
cudaEventDestroy(stop);
printf("DONE in %fs.\n\n", 0.001*elapsedTime);
//<><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><><>
//allocate memory for the resampled image for the return
float *imr = (float*)malloc(Cim.VXNRx*Cim.VXNRy*Cim.VXNRz*sizeof(float));
//copy the image from GPU to CPU
HANDLE_ERROR( cudaMemcpy(imr, d_imr, Cim.VXNRx*Cim.VXNRy*Cim.VXNRz*sizeof(float), cudaMemcpyDeviceToHost) );
return imr;
}
#include <stdio.h>
#ifndef RSMPL_H
#define RSMPL_H
//fine subsampling divisor
#define VDIV 10
#define HANDLE_ERROR(err) (HandleError( err, __FILE__, __LINE__ ))
void HandleError( cudaError_t err, const char *file, int line );
struct Cimg{
float VXSOx;
float VXSOy;
float VXSOz;
short VXNOx;
short VXNOy;
short VXNOz;
float OFFOx;
float OFFOy;
float OFFOz;
float VXSRx;
float VXSRy;
float VXSRz;
short VXNRx;
short VXNRy;
short VXNRz;
float OFFRx;
float OFFRy;
float OFFRz;
};
float *rsmpl( float *imo,
float *A,
Cimg Cim);
#endif
\ No newline at end of file
OBJDIR=obj
SRCDIR=src
INCDIR=..
INCDIR2=../src
CC=nvcc
NFLAGS=-arch=sm_35 --ptxas-options=-v -Xcompiler '-m64','-fPIC','-DNDEBUG','-O3','-Wall','-Wstrict-prototypes'
CFLAGS=-I/usr/include/python2.7 -I$(INCDIR) -I$(INCDIR2)
LDFLAGS=-L/usr/local/lib -L/usr/local/cuda/lib64 -lcudart -lcurand
DEPS = $(INCDIR)/def.h $(INCDIR2)/aux.h $(SRCDIR)/lmaux.h $(SRCDIR)/lmproc.h $(SRCDIR)/hst.h $(SRCDIR)/rnd.h
OBJ = $(OBJDIR)/lmaux.o $(OBJDIR)/lm_module.o $(OBJDIR)/lmproc.o $(OBJDIR)/hst.o $(OBJDIR)/rnd.o
ODIR = ./
$(OBJDIR)/%.o: $(SRCDIR)/%.cu $(DEPS)
$(CC) $(NFLAGS) -c -o $@ $< $(CFLAGS)
$(ODIR)mmr_lmproc.so: $(OBJ)
gcc -shared -o $@ $^ $(LDFLAGS)
$(OBJ): | $(OBJDIR)
$(OBJDIR):
mkdir -p ./$(OBJDIR)
\ No newline at end of file
# init the package folder
import mmrhist
\ No newline at end of file
This diff is collapsed.
This diff is collapsed.
#ifndef HST_H
#define HST_H
#include "aux.h"
#include "lmaux.h"
#include <cuda.h>
#include <curand_kernel.h>
#include <curand.h>
extern LMprop lmprop;
extern int* lm;
curandStatePhilox4_32_10_t* setup_curand();
void gpu_hst(unsigned int *d_ssrb,
unsigned int *d_sino,
unsigned int *d_rdlyd,
unsigned int *d_rprmt,
mMass d_mass,
unsigned int *d_snview,
unsigned int *d_fansums,
unsigned int *d_bucks,
int tstart, int tstop, int nbtp,
LORcc *s2cF,
axialLUT axLUT,
const Cnst Cnt);
#define min(a, b) ({__typeof__(a) _a = (a); \
__typeof__(b) _b = (b); \
_a < _b ? _a : _b;})
#endif
This diff is collapsed.
This diff is collapsed.
#include <stdio.h>
#include "def.h"
#ifndef LAUX_H
#define LAUX_H
#define HANDLE_ERROR(err) (HandleError( err, __FILE__, __LINE__ ))
void HandleError( cudaError_t err, const char *file, int line );
extern LMprop lmprop;
void getMemUse(void);
//get the properties of LM and the chunks into which the LM is divided
void getLMinfo(char *flm);
//modify the properties of LM in case of dynamic studies as the number of frames wont fit in the memory
void modifyLMinfo(int tstart, int tstop);
//uncompress the sinogram after GPU execution
void dsino_ucmpr(unsigned int *d_dsino,
unsigned char *pdsn, unsigned char *ddsn,
int tot_bins, int nfrm);
#endif //LAUX_H
This diff is collapsed.
#ifndef LMPROC_H
#define LMPROC_H
#include <stdlib.h>
#include "def.h"
#include "aux.h"
#include "lmaux.h"
#include "auxmath.h"
#include "hst.h"
typedef struct{
int nitag;
int sne; //number of elements in sino views
unsigned int * snv; //sino views
unsigned int * hcp; //head curve prompts
unsigned int * hcd; //head curve delayeds
unsigned int * fan; //fansums
unsigned int * bck; //buckets (singles)
float * mss; //centre of mass (axially)
unsigned int * ssr;
void * psn;
void * dsn;
unsigned long long psm;
unsigned long long dsm;