Commit 8a5d3e10 authored by breuera's avatar breuera

Added compile options for a compute capability < 2.0.

parent 652d8bf3
...@@ -73,7 +73,8 @@ vars.AddVariables( ...@@ -73,7 +73,8 @@ vars.AddVariables(
), ),
EnumVariable( 'computeCapability', 'optional architecture/compute capability of the CUDA card', 'sm_20', EnumVariable( 'computeCapability', 'optional architecture/compute capability of the CUDA card', 'sm_20',
allowed_values=('sm_20', 'sm_21', 'sm_22', 'sm_23') allowed_values=('sm_10', 'sm_11', 'sm_12','sm_13',
'sm_20', 'sm_21', 'sm_22', 'sm_23' )
), ),
BoolVariable( 'openGL', 'compile with OpenGL visualization', False), BoolVariable( 'openGL', 'compile with OpenGL visualization', False),
......
...@@ -192,9 +192,13 @@ void computeNetUpdatesKernel( ...@@ -192,9 +192,13 @@ void computeNetUpdatesKernel(
reductionPartner = computeOneDPositionKernel(threadIdx.x, threadIdx.y+reductionBlockDimY, blockDim.y); reductionPartner = computeOneDPositionKernel(threadIdx.x, threadIdx.y+reductionBlockDimY, blockDim.y);
} }
#ifndef NDEBUG #ifndef NDEBUG
#if defined(__CUDA_ARCH__) & (__CUDA_ARCH__ < 200)
#warning skipping printf command, which was introduced for compute capability >= 2.0
#else
else { else {
printf("computeNetUpdatesKernel(...): There was an error in the reducing procedure!\n"); printf("computeNetUpdatesKernel(...): There was an error in the reducing procedure!\n");
} }
#endif
#endif #endif
if(threadIdx.x < reductionBlockDimX && threadIdx.y < reductionBlockDimY) { // use only half the threads in each reduction if(threadIdx.x < reductionBlockDimX && threadIdx.y < reductionBlockDimY) { // use only half the threads in each reduction
//execute the reduction routine (maximum) //execute the reduction routine (maximum)
...@@ -307,8 +311,12 @@ void updateUnknownsKernel( ...@@ -307,8 +311,12 @@ void updateUnknownsKernel(
l_cellPosition = computeOneDPositionKernel(l_cellIndexI, l_cellIndexJ, i_nY+2); l_cellPosition = computeOneDPositionKernel(l_cellIndexI, l_cellIndexJ, i_nY+2);
#ifndef NDEBUG #ifndef NDEBUG
#if defined(__CUDA_ARCH__) & (__CUDA_ARCH__ < 200)
#warning skipping printf command, which was introduced for compute capability >= 2.0
#else
if(l_cellPosition > (i_nX+2)*(i_nY+2)) if(l_cellPosition > (i_nX+2)*(i_nY+2))
printf("Warning: cellPosition(%i) > (i_nx+2)*(i_ny+2)\n", l_cellPosition); printf("Warning: cellPosition(%i) > (i_nx+2)*(i_ny+2)\n", l_cellPosition);
#endif
#endif #endif
//! positions of the net-updates in the global CUDA-arrays. //! positions of the net-updates in the global CUDA-arrays.
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment