[gambit-list] embedding CUDA-C (for nVIDIA GPU usage) in Gambit?

Jason E. Aten j.e.aten at gmail.com
Sun Feb 13 19:47:14 EST 2011


I'll try to answer, but feel free to ask for clarification if I don't
illuminate the topic.

nvcc is really just a coordinator for many calls.  nvcc makes a series of
calls, producing intermediate code (PTX files), that can then be
just-in-time compiled into final (.cubin) architecture specific code by the
nvida libraries.  In a sense, yes, it compiles source code, but then breaks
out kernel code to be targeted to the GPU device architecture(s).  Later in
the compilation sequence they all get merged into one file.

I attach an example of a C++ file that includes at the top in the comments
the 15 commands issued by the nvcc nVidia compiler coordinator.  The final
(15th command) is the link stage, producing a final executable that depends
upon several shared libraries.  The matrixMul.cu example from the SDK is
consolidated into one file, matrixMul.cu.txt, and the commands used to
assemble the final binary broken out at the top of the file.

Thanks again.

Best regards,
Jason


-- 
Jason E. Aten, Ph.D.


> Perhaps I don't understand how nvcc and CUDA work.  Is it the case that
> all source code is compiled by nvcc and then linked to create an executable,
> or is there some other preprocessor/compiler/linker/loader that is involved?
> >
> > Marc
> >
> >
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://mailman.iro.umontreal.ca/pipermail/gambit-list/attachments/20110213/67467e09/attachment.htm>
-------------- next part --------------
/* Original location: ~/NVIDIA_GPU_Computing_SDK/C/src/matrixMul/matrixMul.cu
   NB: this code will only work on version >= 2.0 CUDA devices

   jea: begin compilation instructions /breakdown of commands issued, discerned by adding
   NVCCFLAGS       := -v -keep
   in the ../../common/common.mk  makefile addin.

# The sequence of 15 commands to produce an executable for running on a CUDA 2.1 device
#  are documented below. (version 1.0 device commands are also produced, but are omitted
#  by default by the nvidia make files and compiler coordinator nvcc.
#
# INPUT: this file, matrixMul.cu (I integrated the files matrixMul_kernel.cu and matrixMul.h
#        into this file matrixMul.cu to make it a single file as input).
#
# OUTPUT: matrixMul executable binary, that loads the shared nvidia libraries 
#           ( -lcutil_x86_64 -lshrutil_x86_64 -lcuda -lcudart  )


export  _SPACE_= 
export  _CUDART_=cudart
export  _HERE_=/usr/local/cuda/bin
export  _THERE_=/usr/local/cuda/bin
export  _TARGET_SIZE_=64
export  TOP=/usr/local/cuda/bin/..
export  LD_LIBRARY_PATH=/usr/local/cuda/bin/../lib:/usr/local/cuda/bin/../extools/lib:/usr/local/cula/lib64:/usr/lib
export  PATH=/usr/local/cuda/bin/../open64/bin:/usr/local/cuda/bin:/usr/local/bin:/usr/local/cuda/bin:/home/jaten/uns/bin:/usr/local/bin:/usr/local/cuda/bin:/home/jaten/uns/bin:/usr/local/bin:/usr/local/cuda/bin:/home/jaten/uns/bin:/usr/local/bin:/usr/local/cuda/bin:/home/jaten/uns/bin:/usr/local/sbin:/usr/local/bin:/usr/sbin:/usr/bin:/sbin:/bin:/usr/games
export  INCLUDES="-I/usr/local/cuda/bin/../include -I/usr/local/cuda/bin/../include/cudart"  
export  LIBRARIES="-L/usr/local/cuda/bin/../lib64 -lcudart"
export  CUDAFE_FLAGS=
export  OPENCC_FLAGS=
export  PTXAS_FLAGS=

# (0) input: matrixMul.cu         (host code)
#     input: matrixMul_kernel.cu  (device code, where is this read?)
#     input: matrixMul.h          (included by both of the above)

# (1) 1st gcc outputs : matrixMul.compute_20.cpp1.ii

gcc -D__CUDA_ARCH__=200 -E -x c++ -DCUDA_DOUBLE_MATH_FUNCTIONS  "-I/usr/local/cuda/include" "-I/usr/local/cuda/include/cudart"   -I. -D__CUDACC__ -C  -fno-strict-aliasing -O2 -I"." -I"/usr/local/cuda/include" -I"../../common/inc" -I"../../../shared//inc" -D"UNIX" -include "cuda_runtime.h" -m64 -o "matrixMul.compute_20.cpp1.ii" "matrixMul.cu" 



# (2) CUDA Front End : output: matrixMul.compute_20.cudafe1.c matrixMul.compute_20.cudafe1.gpu matrixMul.compute_20.cudafe1.stub.c

cudafe --m64 --gnu_version=40405 -tused --no_remove_unneeded_entities  --gen_c_file_name "matrixMul.compute_20.cudafe1.c" --stub_file_name "matrixMul.compute_20.cudafe1.stub.c" --gen_device_file_name "matrixMul.compute_20.cudafe1.gpu" --include_file_name "matrixMul.fatbin.c" "matrixMul.compute_20.cpp1.ii" 

# (3) output:  matrixMul.compute_20.cpp2.i

gcc -D__CUDA_ARCH__=200 -E -x c -DCUDA_DOUBLE_MATH_FUNCTIONS  "-I/usr/local/cuda/include" "-I/usr/local/cuda/include/cudart"   -I. -D__CUDACC__ -C  -fno-strict-aliasing -O2 -D__CUDA_PREC_DIV -D__CUDA_PREC_SQRT -I"." -I"/usr/local/cuda/include" -I"../../common/inc" -I"../../../shared//inc" -m64 -o "matrixMul.compute_20.cpp2.i" "matrixMul.compute_20.cudafe1.gpu" 



# (4) ouptput:  matrixMul.compute_20.cudafe2.c matrixMul.compute_20.cudafe2.gpu matrixMul.compute_20.cudafe2.stub.c

cudafe --m64 --gnu_version=40405 --c  --gen_c_file_name "matrixMul.compute_20.cudafe2.c" --stub_file_name "matrixMul.compute_20.cudafe2.stub.c" --gen_device_file_name "matrixMul.compute_20.cudafe2.gpu" --include_file_name "matrixMul.fatbin.c" "matrixMul.compute_20.cpp2.i" 



# (5)  output:  matrixMul.compute_20.cpp3.i 

gcc -D__CUDA_ARCH__=200 -E -x c -DCUDA_DOUBLE_MATH_FUNCTIONS  "-I/usr/local/cuda/include" "-I/usr/local/cuda/include/cudart"   -I. -D__CUDABE__  -fno-strict-aliasing -O2 -D__CUDA_PREC_DIV -D__CUDA_PREC_SQRT -I"." -I"/usr/local/cuda/include" -I"../../common/inc" -I"../../../shared//inc" -m64 -o "matrixMul.compute_20.cpp3.i" "matrixMul.compute_20.cudafe2.gpu" 



# (6) output:  matrixMul.hash

filehash -s " " "matrixMul.compute_20.cpp3.i" > "matrixMul.hash"


# (7) output:  matrixMul.cpp4.ii

gcc -E -x c++ "-I/usr/local/cuda/include" "-I/usr/local/cuda/include/cudart"   -I. -D__CUDACC__ -C  -fno-strict-aliasing -O2 -I"." -I"/usr/local/cuda/include" -I"../../common/inc" -I"../../../shared//inc" -D"UNIX" -include "cuda_runtime.h" -m64 -o "matrixMul.cpp4.ii" "matrixMul.cu" 



# (8) output:  matrixMul.compute_20.cudafe1.cpp

cudafe++ --m64 --gnu_version=40405 --parse_templates  --gen_c_file_name "matrixMul.compute_20.cudafe1.cpp" --stub_file_name "matrixMul.compute_20.cudafe1.stub.c" "matrixMul.cpp4.ii" 


# (9) output:  matrixMul.compute_20.ptx

nvopencc  -TARG:compute_20 -m64 -CG:ftz=0 -CG:prec_div=1 -CG:prec_sqrt=1  "matrixMul.compute_20" "matrixMul.compute_20.cpp3.i"  -o "matrixMul.compute_20.ptx"


# (10) output:  matrixMul.compute_20.sm_20.cubin

ptxas  -arch=sm_20 -m64  "matrixMul.compute_20.ptx"  -o "matrixMul.compute_20.sm_20.cubin" 



# (11) output: matrixMul.cpp4.ii  (again!; same as 7)

gcc -E -x c++ "-I/usr/local/cuda/include" "-I/usr/local/cuda/include/cudart"   -I. -D__CUDACC__ -C  -fno-strict-aliasing -O2 -I"." -I"/usr/local/cuda/include" -I"../../common/inc" -I"../../../shared//inc" -D"UNIX" -include "cuda_runtime.h" -m64 -o "matrixMul.cpp4.ii" "matrixMul.cu" 


# (12) output: matrixMul.fatbin.c

fatbin --key="15530d92dba8868a" --source-name="matrixMul.cu" --usage-mode=" " --embedded-fatbin="matrixMul.fatbin.c" "--image=profile=compute_20,file=matrixMul.compute_20.ptx" "--image=profile=sm_20 at compute_20,file=matrixMul.compute_20.sm_20.cubin"


# (13) output: matrixMul.cu.cpp

gcc -D__CUDA_ARCH__=200 -E -x c++ -DCUDA_DOUBLE_MATH_FUNCTIONS  "-I/usr/local/cuda/include" "-I/usr/local/cuda/include/cudart"   -I.  -fno-strict-aliasing -O2 -D__CUDA_PREC_DIV -D__CUDA_PREC_SQRT -I"." -I"/usr/local/cuda/include" -I"../../common/inc" -I"../../../shared//inc" -m64 -o "matrixMul.cu.cpp" "matrixMul.compute_20.cudafe1.cpp" 


# (14) output: obj/x86_64/release/matrixMul.cu.o

gcc -c -x c++ "-I/usr/local/cuda/include" "-I/usr/local/cuda/include/cudart"   -I. -fno-strict-aliasing -O2 -I"." -I"/usr/local/cuda/include" -I"../../common/inc" -I"../../../shared//inc" -fpreprocessed -m64 -o "obj/x86_64/release/matrixMul.cu.o" "matrixMul.cu.cpp" 


# (15) link into ./matrixMul

g++ -m64 -o ./matrixMul  obj/x86_64/release/matrixMul.cu.o  -L/usr/local/cuda/lib64 -L/home/jaten/NVIDIA_GPU_Computing_SDK/C/lib/ -L/home/jaten/NVIDIA_GPU_Computing_SDK/shared/lib/    -lcutil_x86_64 -lshrutil_x86_64 -lcuda -lcudart 




   jea: end compilation breakdown
   jea: being input of matrixMul.cu file (which originally #included matrixMul_kernel.cu and matrixMul.h, but those are now built into this file for simplicity of transport )

*/ 


/*
 * Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.
 *
 * Please refer to the NVIDIA end user license agreement (EULA) associated
 * with this source code for terms and conditions that govern your use of
 * this software. Any use, reproduction, disclosure, or distribution of
 * this software and related documentation outside the terms of the EULA
 * is strictly prohibited.
 *
 */

/* Matrix multiplication: C = A * B.
 * Host code.
 *
 * This sample implements matrix multiplication as described in Chapter 3
 * of the programming guide.
 * It has been written for clarity of exposition to illustrate various CUDA
 * programming principles, not with the goal of providing the most
 * performant generic kernel for matrix multiplication.
 *
 * CUBLAS provides high-performance matrix multiplication.
 * See also:
 * V. Volkov and J. Demmel, "Benchmarking GPUs to tune dense linear algebra,"
 * in Proc. 2008 ACM/IEEE Conf. on Superconducting (SC '08),
 * Piscataway, NJ: IEEE Press, 2008, pp. Art. 31:1-11. 
 *
 */

// Utilities and system includes
#include <shrUtils.h>
#include "cutil_inline.h"

// includes, kernels
//jea put it all in one file for now, instead of: include <matrixMul_kernel.cu>

////////////////////////////////////////
////////////////////////////////////////
//////// begin kernel, from matrixMul_kernel.cu
////////////////////////////////////////
////////////////////////////////////////
/*
 * Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.
 *
 * Please refer to the NVIDIA end user license agreement (EULA) associated
 * with this source code for terms and conditions that govern your use of
 * this software. Any use, reproduction, disclosure, or distribution of
 * this software and related documentation outside the terms of the EULA
 * is strictly prohibited.
 *
 */

/* Matrix multiplication: C = A * B.
 * Device code.
 */

#ifndef _MATRIXMUL_KERNEL_H_
#define _MATRIXMUL_KERNEL_H_

#include <stdio.h>
//// include "matrixMul.h"

//// begin matrixMul.h

#ifndef _MATRIXMUL_H_
#define _MATRIXMUL_H_

// Thread block size
#define BLOCK_SIZE 16

// Basic Matrix dimensions (can be amplified by command line switch)
// (chosen as multiples of the thread block size for simplicity)
#define WA (5  * BLOCK_SIZE) // Matrix A width
#define HA (10 * BLOCK_SIZE) // Matrix A height
#define WB (5  * BLOCK_SIZE) // Matrix B width
#define HB WA  // Matrix B height
#define WC WB  // Matrix C width 
#define HC HA  // Matrix C height

#endif // _MATRIXMUL_H_



/// end matrixMul.h

//// resume matrixMul_kernel.cu

#define CHECK_BANK_CONFLICTS 0
#if CHECK_BANK_CONFLICTS
#define AS(i, j) cutilBankChecker(((float*)&As[0][0]), (BLOCK_SIZE * i + j))
#define BS(i, j) cutilBankChecker(((float*)&Bs[0][0]), (BLOCK_SIZE * i + j))
#else
#define AS(i, j) As[i][j]
#define BS(i, j) Bs[i][j]
#endif

////////////////////////////////////////////////////////////////////////////////
//! Matrix multiplication on the device: C = A * B
//! wA is A's width and wB is B's width
////////////////////////////////////////////////////////////////////////////////
__global__ void
matrixMul( float* C, float* A, float* B, int wA, int wB)
{
    // Block index
    int bx = blockIdx.x;
    int by = blockIdx.y;

    // Thread index
    int tx = threadIdx.x;
    int ty = threadIdx.y;

    // Index of the first sub-matrix of A processed by the block
    int aBegin = wA * BLOCK_SIZE * by;

    // Index of the last sub-matrix of A processed by the block
    int aEnd   = aBegin + wA - 1;

    // Step size used to iterate through the sub-matrices of A
    int aStep  = BLOCK_SIZE;

    // Index of the first sub-matrix of B processed by the block
    int bBegin = BLOCK_SIZE * bx;

    // Step size used to iterate through the sub-matrices of B
    int bStep  = BLOCK_SIZE * wB;

    // Csub is used to store the element of the block sub-matrix
    // that is computed by the thread
    float Csub = 0;

    // Loop over all the sub-matrices of A and B
    // required to compute the block sub-matrix
    for (int a = aBegin, b = bBegin;
             a <= aEnd;
             a += aStep, b += bStep) {

        // Declaration of the shared memory array As used to
        // store the sub-matrix of A
        __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];

        // Declaration of the shared memory array Bs used to
        // store the sub-matrix of B
        __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];

        // Load the matrices from device memory
        // to shared memory; each thread loads
        // one element of each matrix
        AS(ty, tx) = A[a + wA * ty + tx];
        BS(ty, tx) = B[b + wB * ty + tx];

        // Synchronize to make sure the matrices are loaded
        __syncthreads();

        // Multiply the two matrices together;
        // each thread computes one element
        // of the block sub-matrix
        for (int k = 0; k < BLOCK_SIZE; ++k)
            Csub += AS(ty, k) * BS(k, tx);

        // Synchronize to make sure that the preceding
        // computation is done before loading two new
        // sub-matrices of A and B in the next iteration
        __syncthreads();
    }

    // Write the block sub-matrix to device memory;
    // each thread writes one element
    int c = wB * BLOCK_SIZE * by + BLOCK_SIZE * bx;
    C[c + wB * ty + tx] = Csub;
}

#endif // #ifndef _MATRIXMUL_KERNEL_H_


////////////////////////////////////////
////////////////////////////////////////
//////// end kernel, from matrixMul_kernel.cu
////////////////////////////////////////
////////////////////////////////////////

static char *sSDKsample = "matrixMul";

////////////////////////////////////////////////////////////////////////////////
// declaration, forward
void runTest(int argc, char** argv);
void randomInit(float*, int);
void printDiff(float*, float*, int, int, int, float);

extern "C"
void computeGold(float*, const float*, const float*, unsigned int, unsigned int, unsigned int);

// jea: used to be in matrixMul_gold.cpp, but simplify a little by including it here.
void
computeGold(float* C, const float* A, const float* B, unsigned int hA, unsigned int wA, unsigned int wB)
{
    for (unsigned int i = 0; i < hA; ++i)
        for (unsigned int j = 0; j < wB; ++j) {
            double sum = 0;
            for (unsigned int k = 0; k < wA; ++k) {
                double a = A[i * wA + k];
                double b = B[k * wB + j];
                sum += a * b;
            }
            C[i * wB + j] = (float)sum;
        }
}

////////////////////////////////////////////////////////////////////////////////
// Program main
////////////////////////////////////////////////////////////////////////////////
int main(int argc, char** argv)
{
	printf("[ %s ]\n", sSDKsample);

    shrSetLogFileName ("matrixMul.txt");
    shrLog("%s Starting...\n\n", argv[0]);

    runTest(argc, argv);

   // shrEXIT(argc, (const char**)argv);
}

////////////////////////////////////////////////////////////////////////////////
//! Run a simple test for CUDA
////////////////////////////////////////////////////////////////////////////////
void runTest(int argc, char** argv)
{
    if(shrCheckCmdLineFlag(argc, (const char**)argv, "device"))
    {
        cutilDeviceInit(argc, argv);
    }
    else
    {
        cudaSetDevice(cutGetMaxGflopsDeviceId());
    }

    int devID;
    cudaDeviceProp props;

    // get number of SMs on this GPU
    cutilSafeCall(cudaGetDevice(&devID));
    cutilSafeCall(cudaGetDeviceProperties(&props, devID));

    printf("Device %d: \"%s\" with Compute %d.%d capability\n", devID, props.name, props.major, props.minor);

	// set seed for rand()
    srand(2006);

    // Optional Command-line multiplier for matrix sizes
    unsigned int uiWA, uiHA, uiWB, uiHB, uiWC, uiHC;
    int iSizeMultiple = 1;
    shrGetCmdLineArgumenti(argc, (const char**)argv, "sizemult", &iSizeMultiple); 
    iSizeMultiple = CLAMP(iSizeMultiple, 1, 10);

	// For GPUs with fewer # of SM's, we limit the maximum size of the matrix
	if (props.multiProcessorCount <= 4) {
		uiWA = 2 * BLOCK_SIZE * iSizeMultiple;
		uiHA = 4 * BLOCK_SIZE * iSizeMultiple;
		uiWB = 2 * BLOCK_SIZE * iSizeMultiple;
		uiHB = 4 * BLOCK_SIZE * iSizeMultiple;
		uiWC = 2 * BLOCK_SIZE * iSizeMultiple;
		uiHC = 4 * BLOCK_SIZE * iSizeMultiple;
	} else {
		uiWA = WA * iSizeMultiple;
		uiHA = HA * iSizeMultiple;
		uiWB = WB * iSizeMultiple;
		uiHB = HB * iSizeMultiple;
		uiWC = WC * iSizeMultiple;
		uiHC = HC * iSizeMultiple;
	}
    shrLog("\nUsing Matrix Sizes: A(%u x %u), B(%u x %u), C(%u x %u)\n\n", 
            uiWA, uiHA, uiWB, uiHB, uiWC, uiHC);

    // allocate host memory for matrices A and B
    unsigned int size_A = uiWA * uiHA;
    unsigned int mem_size_A = sizeof(float) * size_A;
    float* h_A = (float*)malloc(mem_size_A);
    unsigned int size_B = uiWB * uiHB;
    unsigned int mem_size_B = sizeof(float) * size_B;
    float* h_B = (float*)malloc(mem_size_B);

    // initialize host memory
    randomInit(h_A, size_A);
    randomInit(h_B, size_B);

    // allocate device memory
    float* d_A;
    cutilSafeCall(cudaMalloc((void**) &d_A, mem_size_A));
    float* d_B;
    cutilSafeCall(cudaMalloc((void**) &d_B, mem_size_B));

    // copy host memory to device
    cutilSafeCall(cudaMemcpy(d_A, h_A, mem_size_A,
                              cudaMemcpyHostToDevice) );
    cutilSafeCall(cudaMemcpy(d_B, h_B, mem_size_B,
                              cudaMemcpyHostToDevice) );

    // allocate device memory for result
    unsigned int size_C = uiWC * uiHC;
    unsigned int mem_size_C = sizeof(float) * size_C;
    float* d_C;
    cutilSafeCall(cudaMalloc((void**) &d_C, mem_size_C));

    // allocate host memory for the result
    float* h_C = (float*) malloc(mem_size_C);
    

    // setup execution parameters
    dim3 threads(BLOCK_SIZE, BLOCK_SIZE);
    dim3 grid(uiWC / threads.x, uiHC / threads.y);

    // kernel warmup
    matrixMul<<< grid, threads >>>(d_C, d_A, d_B, uiWA, uiWB);
    cudaThreadSynchronize();
    
    // create and start timer
    shrLog("Run Kernels...\n\n");
    unsigned int timer = 0;
    cutilCheckError(cutCreateTimer(&timer));
    cutilCheckError(cutStartTimer(timer));

    // execute the kernel
    int nIter = 30;
    for (int j = 0; j < nIter; j++) 
		{
            matrixMul<<< grid, threads >>>(d_C, d_A, d_B, uiWA, uiWB);
        }

    // check if kernel execution generated and error
    cutilCheckMsg("Kernel execution failed");

    cudaThreadSynchronize();
    // stop and destroy timer
    cutilCheckError(cutStopTimer(timer));
    double dSeconds = cutGetTimerValue(timer)/((double)nIter * 1000.0);
    double dNumOps = 2.0 * (double)uiWA * (double)uiHA * (double)uiWB;
    double gflops = 1.0e-9 * dNumOps/dSeconds;

    //Log througput, etc
    shrLogEx(LOGBOTH | MASTER, 0, "matrixMul, Throughput = %.4f GFlop/s, Time = %.5f s, Size = %.0f Ops, NumDevsUsed = %d, Workgroup = %u\n", 
            gflops, dSeconds, dNumOps, 1, threads.x * threads.y);
    cutilCheckError(cutDeleteTimer(timer));

    // copy result from device to host
    cutilSafeCall(cudaMemcpy(h_C, d_C, mem_size_C,
                              cudaMemcpyDeviceToHost) );

    // compute reference solution
    shrLog("\nCheck against Host computation...\n\n");    
    float* reference = (float*)malloc(mem_size_C);
    computeGold(reference, h_A, h_B, uiHA, uiWA, uiWB);

    // check result
    shrBOOL res = shrCompareL2fe(reference, h_C, size_C, 1.0e-6f);
    if (res != shrTRUE) 
    {
        printDiff(reference, h_C, uiWC, uiHC, 100, 1.0e-5f);
    }
    shrLog("%s \n\n", (shrTRUE == res) ? "PASSED" : "FAILED");

    // clean up memory
    free(h_A);
    free(h_B);
    free(h_C);
    free(reference);
    cutilSafeCall(cudaFree(d_A));
    cutilSafeCall(cudaFree(d_B));
    cutilSafeCall(cudaFree(d_C));

    cudaThreadExit();
}

// Allocates a matrix with random float entries.
void randomInit(float* data, int size)
{
    for (int i = 0; i < size; ++i)
        data[i] = rand() / (float)RAND_MAX;
}

void printDiff(float *data1, float *data2, int width, int height, int iListLength, float fListTol)
{
    shrLog("Listing first %d Differences > %.6f...\n", iListLength, fListTol);
    int i,j,k;
    int error_count=0;
    for (j = 0; j < height; j++) 
    {
        if (error_count < iListLength)
        {
            shrLog("\n  Row %d:\n", j);
        }
        for (i = 0; i < width; i++) 
        {
            k = j * width + i;
            float fDiff = fabs(data1[k] - data2[k]);
            if (fDiff > fListTol) 
            {                
                if (error_count < iListLength)
                {
                    shrLog("    Loc(%d,%d)\tCPU=%.5f\tGPU=%.5f\tDiff=%.6f\n", i, j, data1[k], data2[k], fDiff);
                }
                error_count++;
            }
        }
    }
    shrLog(" \n  Total Errors = %d\n\n", error_count);
}


More information about the Gambit-list mailing list