void CudaModule::launchKernel(CUfunction kernel, const Vec2i& blockSize,
const Vec2i& gridSize, bool async,
CUstream stream)
{
if (!kernel) {
fail("CudaModule: No kernel specified!");
}
#if (CUDA_VERSION >= 3000)
if (NULL != cuFuncSetCacheConfig)
{
CUfunc_cache cache = (s_preferL1)? CU_FUNC_CACHE_PREFER_L1 :
CU_FUNC_CACHE_PREFER_SHARED;
checkError("cuFuncSetCacheConfig", cuFuncSetCacheConfig( kernel, cache) );
}
#endif
updateGlobals();
updateTexRefs(kernel);
checkError("cuFuncSetBlockShape", cuFuncSetBlockShape(kernel, blockSize.x, blockSize.y, 1));
if (async && (NULL != cuLaunchGridAsync))
{
checkError("cuLaunchGridAsync",
cuLaunchGridAsync(kernel, gridSize.x, gridSize.y, stream));
}
else
{
checkError("cuLaunchGrid",
cuLaunchGrid(kernel, gridSize.x, gridSize.y));
}
}
const unsigned long CUDARunner::RunStep()
{
//unsigned int best=0;
//unsigned int bestg=~0;
int offset=0;
if(m_in==0 || m_out==0 || m_devin==0 || m_devout==0)
{
AllocateResources(m_numb,m_numt);
}
m_out[0].m_bestnonce=0;
cuMemcpyHtoD(m_devout,m_out,/*m_numb*m_numt*/sizeof(cuda_out));
cuMemcpyHtoD(m_devin,m_in,sizeof(cuda_in));
int loops=GetStepIterations();
int bits=GetStepBitShift()-1;
void *ptr=(void *)(size_t)m_devin;
ALIGN_UP(offset, __alignof(ptr));
cuParamSetv(m_function,offset,&ptr,sizeof(ptr));
offset+=sizeof(ptr);
ptr=(void *)(size_t)m_devout;
ALIGN_UP(offset, __alignof(ptr));
cuParamSetv(m_function,offset,&ptr,sizeof(ptr));
offset+=sizeof(ptr);
ALIGN_UP(offset, __alignof(loops));
cuParamSeti(m_function,offset,loops);
offset+=sizeof(loops);
ALIGN_UP(offset, __alignof(bits));
cuParamSeti(m_function,offset,bits);
offset+=sizeof(bits);
cuParamSetSize(m_function,offset);
cuFuncSetBlockShape(m_function,m_numt,1,1);
cuLaunchGrid(m_function,m_numb,1);
cuMemcpyDtoH(m_out,m_devout,/*m_numb*m_numt*/sizeof(cuda_out));
// very unlikely that we will find more than 1 hash with H=0
// so we'll just return the first one and not even worry about G
for(int i=0; i<1/*m_numb*m_numt*/; i++)
{
if(m_out[i].m_bestnonce!=0)// && m_out[i].m_bestg<bestg)
{
return CryptoPP::ByteReverse(m_out[i].m_bestnonce);
//best=m_out[i].m_bestnonce;
//bestg=m_out[i].m_bestg;
}
}
return 0;
}
/*
* Class: edu_syr_pcpratts_rootbeer_runtime2_cuda_CudaRuntime2
* Method: runBlocks
* Signature: (I)V
*/
JNIEXPORT jint JNICALL Java_edu_syr_pcpratts_rootbeer_runtime2_cuda_CudaRuntime2_runBlocks
(JNIEnv *env, jobject this_obj, jint num_blocks, jint block_shape, jint grid_shape){
CUresult status;
jlong * infoSpace = (jlong *) malloc(gc_space_size);
infoSpace[1] = heapEndPtr;
cuMemcpyHtoD(gcInfoSpace, infoSpace, gc_space_size);
cuMemcpyHtoD(gpuToSpace, toSpace, heapEndPtr);
//cuMemcpyHtoD(gpuTexture, textureMemory, textureMemSize);
cuMemcpyHtoD(gpuHandlesMemory, handlesMemory, num_blocks * sizeof(jlong));
cuMemcpyHtoD(gpuHeapEndPtr, &heapEndPtr, sizeof(jlong));
cuMemcpyHtoD(gpuBufferSize, &bufferSize, sizeof(jlong));
/*
status = cuModuleGetTexRef(&cache, cuModule, "m_Cache");
if (CUDA_SUCCESS != status)
{
printf("error in cuModuleGetTexRef %d\n", status);
}
status = cuTexRefSetAddress(0, cache, gpuTexture, textureMemSize);
if (CUDA_SUCCESS != status)
{
printf("error in cuTextRefSetAddress %d\n", status);
}
*/
status = cuFuncSetBlockShape(cuFunction, block_shape, 1, 1);
if (CUDA_SUCCESS != status)
{
printf("error in cuFuncSetBlockShape %d\n", status);
return (jint) status;
}
status = cuLaunchGrid(cuFunction, grid_shape, 1);
if (CUDA_SUCCESS != status)
{
printf("error in cuLaunchGrid %d\n", status);
fflush(stdout);
return (jint) status;
}
status = cuCtxSynchronize();
if (CUDA_SUCCESS != status)
{
printf("error in cuCtxSynchronize %d\n", status);
return (jint) status;
}
cuMemcpyDtoH(infoSpace, gcInfoSpace, gc_space_size);
heapEndPtr = infoSpace[1];
cuMemcpyDtoH(toSpace, gpuToSpace, heapEndPtr);
cuMemcpyDtoH(exceptionsMemory, gpuExceptionsMemory, num_blocks * sizeof(jlong));
free(infoSpace);
return 0;
}
//----------------------------------------------------------------------------//
bool CUDAImpl::_LaunchKernel(Kernel & kernel,
const CUfunction & cudaKernel,
std::string * err)
{
// Set CUDA kernel arguments
CUresult c_err;
int paramOffset = 0;
for(size_t i = 0; i < kernel.inBuffers.size(); ++i) {
c_err = cuParamSetv(cudaKernel, paramOffset,
&_cudaBuffers[kernel.inBuffers[i].buffer->name],
sizeof(void*));
paramOffset += sizeof(void *);
}
for(size_t i = 0; i < kernel.outBuffers.size(); ++i) {
c_err = cuParamSetv(cudaKernel, paramOffset,
&_cudaBuffers[kernel.outBuffers[i].buffer->name],
sizeof(void*));
paramOffset += sizeof(void *);
}
for(size_t i = 0; i < kernel.paramsInt.size(); ++i) {
c_err = cuParamSetv(cudaKernel, paramOffset,
&kernel.paramsInt[i].value, sizeof(int));
paramOffset += sizeof(int);
}
for(size_t i = 0; i < kernel.paramsFloat.size(); ++i) {
c_err = cuParamSetv(cudaKernel, paramOffset,
&kernel.paramsFloat[i].value, sizeof(float));
paramOffset += sizeof(float);
}
// int and width parameters
c_err = cuParamSetv(cudaKernel, paramOffset, &_w, sizeof(int));
paramOffset += sizeof(int);
c_err = cuParamSetv(cudaKernel, paramOffset, &_h, sizeof(int));
paramOffset += sizeof(int);
// It should be fine to check once all the arguments have been set
if(_cudaErrorCheckParamSet(c_err, err, kernel.name)) {
return false;
}
c_err = cuParamSetSize(cudaKernel, paramOffset);
if (_cudaErrorParamSetSize(c_err, err, kernel.name)) {
return false;
}
// Launch the CUDA kernel
const int nBlocksHor = _w / 16 + 1;
const int nBlocksVer = _h / 16 + 1;
cuFuncSetBlockShape(cudaKernel, 16, 16, 1);
c_err = cuLaunchGrid(cudaKernel, nBlocksHor, nBlocksVer);
if (_cudaErrorLaunchKernel(c_err, err, kernel.name)) {
return false;
}
return true;
}
void GPUInterface::LaunchKernel(GPUFunction deviceFunction,
Dim3Int block,
Dim3Int grid,
int parameterCountV,
int totalParameterCount,
...) { // parameters
#ifdef BEAGLE_DEBUG_FLOW
fprintf(stderr,"\t\t\tEntering GPUInterface::LaunchKernel\n");
#endif
SAFE_CUDA(cuCtxPushCurrent(cudaContext));
SAFE_CUDA(cuFuncSetBlockShape(deviceFunction, block.x, block.y, block.z));
int offset = 0;
va_list parameters;
va_start(parameters, totalParameterCount);
for(int i = 0; i < parameterCountV; i++) {
void* param = (void*)(size_t)va_arg(parameters, GPUPtr);
// adjust offset alignment requirements
offset = (offset + __alignof(param) - 1) & ~(__alignof(param) - 1);
SAFE_CUDA(cuParamSetv(deviceFunction, offset, ¶m, sizeof(param)));
offset += sizeof(void*);
}
for(int i = parameterCountV; i < totalParameterCount; i++) {
unsigned int param = va_arg(parameters, unsigned int);
// adjust offset alignment requirements
offset = (offset + __alignof(param) - 1) & ~(__alignof(param) - 1);
SAFE_CUDA(cuParamSeti(deviceFunction, offset, param));
offset += sizeof(param);
}
va_end(parameters);
SAFE_CUDA(cuParamSetSize(deviceFunction, offset));
SAFE_CUDA(cuLaunchGrid(deviceFunction, grid.x, grid.y));
SAFE_CUDA(cuCtxPopCurrent(&cudaContext));
#ifdef BEAGLE_DEBUG_FLOW
fprintf(stderr,"\t\t\tLeaving GPUInterface::LaunchKernel\n");
#endif
}
CUresult loadAndRunDualTestFunction(CUmodule *phModule, std::string name, CUdeviceptr &d_data0,
CUdeviceptr &d_data1,
DataStruct *h_data0,
DataStruct *h_data1,
unsigned int memSize,
int thread_x=1,int thread_y=1,int thread_z=1,
int block_x=1, int block_y=1, int block_z=1)
{
// std::cout << " Start Loading" << std::endl;
// load data the to device
cuMemcpyHtoD(d_data0, h_data0, memSize);
cuMemcpyHtoD(d_data1, h_data1, memSize);
// Locate the kernel entry point
CUfunction phKernel = 0;
CUresult status = cuModuleGetFunction(&phKernel, *phModule, name.data());
if (status != CUDA_SUCCESS)
{printf("ERROR: could not load function\n");}
// Set the kernel parameters
status = cuFuncSetBlockShape(phKernel, thread_x, thread_y, thread_z);
if (status != CUDA_SUCCESS)
{printf("ERROR: during setBlockShape\n");}
int paramOffset = 0, size=0;
size = sizeof(CUdeviceptr);
status = cuParamSetv(phKernel, paramOffset, &d_data0, size);
paramOffset += size;
status = cuParamSetv(phKernel, paramOffset, &d_data1, size);
paramOffset += size;
status = cuParamSetSize(phKernel, paramOffset);
if (status != CUDA_SUCCESS)
{printf("ERROR: during cuParamSetv\n");}
// Launch the kernel
status = cuLaunchGrid(phKernel, block_x, block_y);
if (status != CUDA_SUCCESS)
{printf("ERROR: during grid launch\n");}
// std::cout << " launched CUDA kernel!!" << std::endl;
// Copy the result back to the host
status = cuMemcpyDtoH(h_data0, d_data0, memSize);
status = cuMemcpyDtoH(h_data1, d_data1, memSize);
if (status != CUDA_SUCCESS)
{printf("ERROR: during MemcpyDtoH\n");}
}
/**
* Invokes the kernel @f on a @gridDimX x @gridDimY x @gridDimZ grid of blocks.
* Each block contains @blockDimX x @blockDimY x @blockDimZ threads.
* @sharedMemBytes sets the amount of dynamic shared memory that will be
* available to each thread block.
*
* cuLaunchKernel() can optionally be associated to a stream by passing a
* non-zero hStream argument.
*
* Kernel parameters to @f can be specified in one of two ways:
*
* 1) Kernel parameters can be specified via kernelParams. If f has N
* parameters, then kernelParams needs to be an array of N pointers. Each of
* kernelParams[0] through kernelParams[N-1] must point to a region of memory
* from which the actual kernel parameter will be copied. The number of kernel
* parameters and their offsets and sizes do not need to be specified as that
* information is retrieved directly from the kernel's image.
*
* 2) Kernel parameters can also be packaged by the application into a single
* buffer that is passed in via the extra parameter. This places the burden on
* the application of knowing each kernel parameter's size and alignment/
* padding within the buffer. Here is an example of using the extra parameter
* in this manner:
*
* size_t argBufferSize;
* char argBuffer[256];
*
* // populate argBuffer and argBufferSize
*
* void *config[] = {
* CU_LAUNCH_PARAM_BUFFER_POINTER, argBuffer,
* CU_LAUNCH_PARAM_BUFFER_SIZE, &argBufferSize,
* CU_LAUNCH_PARAM_END
* };
* status = cuLaunchKernel(f, gx, gy, gz, bx, by, bz, sh, s, NULL, config);
*
* The extra parameter exists to allow cuLaunchKernel to take additional less
* commonly used arguments. extra specifies a list of names of extra settings
* and their corresponding values. Each extra setting name is immediately
* followed by the corresponding value. The list must be terminated with
* either NULL or CU_LAUNCH_PARAM_END.
*
* CU_LAUNCH_PARAM_END, which indicates the end of the extra array;
* CU_LAUNCH_PARAM_BUFFER_POINTER, which specifies that the next value in
* extra will be a pointer to a buffer containing all the kernel parameters
* for launching kernel f;
* CU_LAUNCH_PARAM_BUFFER_SIZE, which specifies that the next value in extra
* will be a pointer to a size_t containing the size of the buffer specified
* with CU_LAUNCH_PARAM_BUFFER_POINTER;
*
* The error CUDA_ERROR_INVALID_VALUE will be returned if kernel parameters
* are specified with both kernelParams and extra (i.e. both kernelParams and
* extra are non-NULL).
*
* Calling cuLaunchKernel() sets persistent function state that is the same as
* function state set through the following deprecated APIs:
*
* cuFuncSetBlockShape() cuFuncSetSharedSize() cuParamSetSize() cuParamSeti()
* cuParamSetf() cuParamSetv()
*
* When the kernel @f is launched via cuLaunchKernel(), the previous block
* shape, shared size and parameter info associated with @f is overwritten.
*
* Note that to use cuLaunchKernel(), the kernel @f must either have been
* compiled with toolchain version 3.2 or later so that it will contain kernel
* parameter information, or have no kernel parameters. If either of these
* conditions is not met, then cuLaunchKernel() will return
* CUDA_ERROR_INVALID_IMAGE.
*
* Parameters:
* f - Kernel to launch
* gridDimX - Width of grid in blocks
* gridDimY - Height of grid in blocks
* gridDimZ - Depth of grid in blocks
* blockDimX - X dimension of each thread block
* blockDimY - Y dimension of each thread block
* blockDimZ - Z dimension of each thread block
* sharedMemBytes - Dynamic shared-memory size per thread block in bytes
* hStream - Stream identifier
* kernelParams - Array of pointers to kernel parameters
* extra - Extra options
*
* Returns:
* CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
* CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_HANDLE,
* CUDA_ERROR_INVALID_IMAGE, CUDA_ERROR_INVALID_VALUE, CUDA_ERROR_LAUNCH_FAILED,
* CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES, CUDA_ERROR_LAUNCH_TIMEOUT,
* CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING,
* CUDA_ERROR_SHARED_OBJECT_INIT_FAILED
*/
CUresult cuLaunchKernel
(CUfunction f,
unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ,
unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ,
unsigned int sharedMemBytes, CUstream hStream,
void **kernelParams, void **extra)
{
struct gdev_cuda_raw_func *rf;
CUresult res;
int i;
if (hStream) {
GDEV_PRINT("cuLaunchKernel: Stream is not supported.\n");
return CUDA_ERROR_INVALID_HANDLE;
}
if (extra) {
GDEV_PRINT("cuLaunchKernel: Extra Parameters are not supported.\n");
return CUDA_ERROR_INVALID_HANDLE;
}
res = cuFuncSetSharedSize(f, sharedMemBytes);
if (res != CUDA_SUCCESS)
return res;
res = cuFuncSetBlockShape(f, blockDimX, blockDimY, blockDimZ);
if (res != CUDA_SUCCESS)
return res;
rf = &f->raw_func;
for (i = 0; i < rf->param_count; i++) {
void *p = kernelParams[i];
int offset = rf->param_info[i].offset;
uint32_t size = rf->param_info[i].size;
cuParamSetv(f, offset, p, size);
}
res = cuParamSetSize(f, rf->param_size);
if (res != CUDA_SUCCESS)
return res;
res = cuLaunchGrid(f, gridDimX, gridDimY);
if (res != CUDA_SUCCESS)
return res;
return CUDA_SUCCESS;
}
/*
* execute kernel function
*/
void
execute_cuda(void){
res = cuLaunchGrid(function, 1, block_num);
if(res != CUDA_SUCCESS){
printf("cuLaunchGrid() failed: res = %s\n", conv(res));
exit(1);
}
res = cuCtxSynchronize();
if(res != CUDA_SUCCESS){
printf("cuCtxSynchronize() failed: res = %s\n", conv(res));
exit(1);
}
}
int main(int argc, char ** argv)
{
int dev_count = 0;
CUdevice device;
CUcontext context;
CUmodule module;
CUfunction function;
cuInit(0);
cuDeviceGetCount(&dev_count);
if (dev_count < 1) return -1;
cuDeviceGet( &device, 0 );
cuCtxCreate( &context, 0, device );
cuModuleLoad( &module, "hello.cuda_runtime.ptx" );
cuModuleGetFunction( &function, module, "_Z6kernelPf" );
int N = 512;
CUdeviceptr pData;
cuMemAlloc( &pData, N * sizeof(float) );
cuFuncSetBlockShape( function, N, 1, 1 );
cuParamSeti( function, 0, pData );
cuParamSetSize( function, 4 );
cuLaunchGrid( function, 1, 1 );
float * pHostData = new float[N];
cuMemcpyDtoH( pHostData, pData, N * sizeof( float) );
cuMemFree( pData );
delete [] pHostData;
return 0;
}
int main( int argc, char** argv)
{
uint num_threads;
uint num_blocks, block_size;
uint length;
uint nBytes;
int *list;
int status, verbose, c, i, j, logBlocks;
int read_stdin;
struct timeval start_time, end_time;
unsigned long total_time;
CUdevice hDevice;
CUcontext hContext;
CUmodule hModule;
CUfunction bitonicBlockFn;
CUfunction mergeBlocksFn;
CUdeviceptr pDeviceArrayA;
CUdeviceptr pDeviceArrayB;
status = SUCCESS;
verbose = 0;
read_stdin = FALSE;
length = 0;
while ((c = getopt (argc, argv, "dip:vO")) != -1) {
switch (c) {
case 'd':
verbose |= GROSS_DEBUG;
break;
case 'i':
read_stdin = TRUE;
case 'O':
verbose |= OUTPUT;
break;
case 'p':
length = 1 << atoi(optarg);
break;
case 'v':
verbose |= DEBUG;
break;
case '?':
default:
print_usage();
return FAILURE;
}
}
if ( read_stdin == TRUE ) {
/* Read sequence of integers from stdin */
list = (int*) malloc (INIT_INPUT_SIZE * sizeof(int) );
length = readIntegers(list, INIT_INPUT_SIZE);
} else if ( length > 0 ) {
list = (int*) malloc (length * sizeof(int) );
randomInts(list, length);
} else if (optind >= argc) { /* No size was given */
print_usage();
return FAILURE;
} else {
/* Generate our own integers */
length = atoi(argv[optind]);
list = (int*) malloc (length * sizeof(int) );
randomInts(list, length);
}
/*
* Phase 1:
* There will be one thread for each element to be sorted. Each
* block will perform bitonic sort on MAX_THREADS_PER_BLOCK elements.
*/
/* Initialize sizes */
num_threads = _min(length, MAX_THREADS_PER_BLOCK );
num_blocks = (length-1) / MAX_THREADS_PER_BLOCK + 1;
nBytes = length * sizeof(int);
if (verbose & DEBUG) printf("Initializing GPU.\n");
/* Start timing */
gettimeofday(&start_time, NULL);
/* Initialize GPU */
cutilDrvSafeCall( cuInit(0) );
cutilDrvSafeCall( cuDeviceGet(&hDevice, 0) );
cutilDrvSafeCall( cuCtxCreate(&hContext, 0, hDevice) );
cutilDrvSafeCall( cuModuleLoad(&hModule, MODULE_FILE) );
cutilDrvSafeCall( cuModuleGetFunction(&bitonicBlockFn, hModule, BITONIC_BLOCK_FN) );
/* Allocate memory on the device */
cutilDrvSafeCall( cuMemAlloc(&pDeviceArrayA, nBytes) );
cutilDrvSafeCall( cuMemAlloc(&pDeviceArrayB, nBytes) );
cutilDrvSafeCall( cuMemcpyHtoD(pDeviceArrayA, list, nBytes) );
cutilDrvSafeCall( cuFuncSetBlockShape(bitonicBlockFn, num_threads, 1, 1));
cutilDrvSafeCall( cuParamSeti(bitonicBlockFn, 0, pDeviceArrayA) );
cutilDrvSafeCall( cuParamSetSize(bitonicBlockFn, 4) );
/* Execute the kernel on the GPU */
if ( verbose & DEBUG ) printf("Launching bitonic sort kernel with %d blocks and %d threads per block.\n", num_blocks, num_threads);
cutilDrvSafeCall( cuLaunchGrid(bitonicBlockFn, num_blocks, 1) );
/*
* Phase 2:
* At this point each block is a sorted list. Now it's time to merge them.
*/
/* TODO This should go away after development */
if ( verbose & GROSS_DEBUG ) {
cuMemcpyDtoH(list, pDeviceArrayA, nBytes);
for (i=0; i<num_blocks; ++i) {
printf("### Block %d:\n", i);
for (j=0; j<num_threads; ++j) {
printf("%d\n", list[i*num_threads + j]);
}
}
}
i=0;
/* Do we need to merge blocks? */
if ( num_blocks > 1 ) {
/* There will be Log_2(num_blocks) merge steps. */
logBlocks = 0;
for (i=1; i<num_blocks; i *= 2) ++logBlocks;
if ( verbose & DEBUG ) printf("There will be %d merge steps.\n", logBlocks);
block_size = num_threads; /* How big the blocks were in the last grid launch. */
num_threads = num_blocks >> 1; /* Start with blocks/2 threads */
num_blocks = (num_threads-1) / MAX_THREADS_PER_BLOCK + 1;
cutilDrvSafeCall( cuModuleGetFunction(&mergeBlocksFn, hModule, MERGE_BLOCKS_FN) );
cuParamSeti(mergeBlocksFn, 4, block_size);
cuParamSetSize(mergeBlocksFn, 16);
for (i=0; i < logBlocks; ++i) {
cuFuncSetBlockShape(mergeBlocksFn, num_threads, 1, 1);
cuParamSeti(mergeBlocksFn, 0, i); /* set merge level */
/* Merging uses a source array and destination array, the gpu has 2 arrays allocated
* so we swap which is the source and which is the destination for each iteration. */
if ( i%2 == 0 ) {
cuParamSeti(mergeBlocksFn, 8, pDeviceArrayA);
cuParamSeti(mergeBlocksFn, 12, pDeviceArrayB);
} else {
cuParamSeti(mergeBlocksFn, 8, pDeviceArrayB);
cuParamSeti(mergeBlocksFn, 12, pDeviceArrayA);
}
if ( verbose & DEBUG ) {
printf("Launching block merge kernel with %d blocks and %d threads per block\n",
num_blocks, num_threads/num_blocks);
}
cutilDrvSafeCall( cuLaunchGrid(mergeBlocksFn, num_blocks, 1) );
num_threads = num_threads >> 1;
num_blocks = (num_threads-1) / MAX_THREADS_PER_BLOCK + 1;
}
}
int main(int argc, char *argv[])
{
argc--; argv++;
// Instruction-level test of PTX assembly language and emulator.
// This test should work natively and under emulation. Many of the
// instructions tested here stress many poorly documented features
// of the PTX assembly language. If the emulator passes these
// tests, then it can surely pass code that is generated by the
// nvcc compiler.
test(cuInit(0), "cuInit");
int deviceCount = 0;
test(cuDeviceGetCount(&deviceCount), "cuDeviceGetCount");
int device = 0;
if (argc)
device = atoi(*argv);
CUdevice cuDevice = 0;
test(cuDeviceGet(&cuDevice, device), "cuDeviceGet");
CUcontext cuContext;
int xxx = cuCtxCreate(&cuContext, 0, cuDevice);
CUmodule cuModule;
test(cuModuleLoad(&cuModule, "inst.ptx"), "cuModuleLoad");
// Do basic test. No sense continuing if we cannot complete this
// test.
try
{
CUfunction proc;
test(cuModuleGetFunction(&proc, cuModule, "InstBasic"), "cuModuleGetFunction");
bool * h_R = (bool*)malloc(sizeof(bool));
memset(h_R, 0, sizeof(bool));
CUdeviceptr d_R;
test(cuMemAlloc(&d_R, sizeof(bool)), "cuMemAlloc");
test(cuMemcpyHtoD(d_R, h_R, sizeof(bool)), "cuMemcpyHtoD");
int offset = 0;
void* ptr;
ptr = (void*)(size_t)d_R;
ALIGN_UP(offset, __alignof(ptr));
test(cuParamSetv(proc, offset, &ptr, sizeof(ptr)), "cuParamSetv");
offset += sizeof(ptr);
test(cuParamSetSize(proc, offset), "cuParamSetSize");
int threadsPerBlock = 1;
int blocksPerGrid = 1;
test(cuFuncSetBlockShape(proc, threadsPerBlock, 1, 1), "cuFuncSetBlockShape");
test(cuLaunchGrid(proc, blocksPerGrid, 1), "cuLaunchGrid");
test(cuMemcpyDtoH(h_R, d_R, sizeof(bool)), "cuMemcpyDtoH");
test(cuMemFree(d_R), "cuMemFree");
if (h_R[0] == 1)
std::cout << "Basic test passed.\n";
else {
std::cout << "Basic test failed.\n";
exit(1);
}
} catch (...)
{
test(1, "test crashed.");
}
// Do LD, ST, MOV test.
load_and_test(cuModule, "InstLSMC");
// Do ADD, SUB test.
load_and_test(cuModule, "InstAddSub");
return 0;
}
int cuda_test_madd_vmmap_hybrid(unsigned int n, char *path)
{
int i, j, idx;
CUresult res;
CUdevice dev;
CUcontext ctx;
CUfunction function;
CUmodule module;
CUdeviceptr a_dev, b_dev, c_dev;
unsigned int *a_buf, *b_buf, *c_buf;
unsigned long long int a_phys, b_phys, c_phys;
unsigned int *c = (unsigned int *) malloc (n*n * sizeof(unsigned int));
int block_x, block_y, grid_x, grid_y;
char fname[256];
int ret = 0;
struct timeval tv;
struct timeval tv_total_start, tv_total_end;
float total;
struct timeval tv_h2d_start, tv_h2d_end;
float h2d;
struct timeval tv_d2h_start, tv_d2h_end;
float d2h;
struct timeval tv_exec_start, tv_exec_end;
struct timeval tv_mem_alloc_start;
struct timeval tv_data_init_start;
float data_init;
struct timeval tv_conf_kern_start;
struct timeval tv_close_start;
float mem_alloc;
float exec;
float init_gpu;
float configure_kernel;
float close_gpu;
float data_read;
unsigned int dummy_b, dummy_c;
/* block_x * block_y should not exceed 512. */
block_x = n < 16 ? n : 16;
block_y = n < 16 ? n : 16;
grid_x = n / block_x;
if (n % block_x != 0)
grid_x++;
grid_y = n / block_y;
if (n % block_y != 0)
grid_y++;
gettimeofday(&tv_total_start, NULL);
res = cuInit(0);
if (res != CUDA_SUCCESS) {
printf("cuInit failed: res = %lu\n", (unsigned long)res);
return -1;
}
res = cuDeviceGet(&dev, 0);
if (res != CUDA_SUCCESS) {
printf("cuDeviceGet failed: res = %lu\n", (unsigned long)res);
return -1;
}
res = cuCtxCreate(&ctx, 0, dev);
if (res != CUDA_SUCCESS) {
printf("cuCtxCreate failed: res = %lu\n", (unsigned long)res);
return -1;
}
sprintf(fname, "%s/madd_gpu.cubin", path);
res = cuModuleLoad(&module, fname);
if (res != CUDA_SUCCESS) {
printf("cuModuleLoad() failed\n");
return -1;
}
res = cuModuleGetFunction(&function, module, "_Z3addPjS_S_j");
if (res != CUDA_SUCCESS) {
printf("cuModuleGetFunction() failed\n");
return -1;
}
res = cuFuncSetBlockShape(function, block_x, block_y, 1);
if (res != CUDA_SUCCESS) {
printf("cuFuncSetBlockShape() failed\n");
return -1;
}
gettimeofday(&tv_mem_alloc_start, NULL);
/* a[] */
res = cuMemAlloc(&a_dev, n*n * sizeof(unsigned int));
if (res != CUDA_SUCCESS) {
printf("cuMemAlloc (a) failed\n");
return -1;
}
res = cuMemMap((void**)&a_buf, a_dev, n*n * sizeof(unsigned int));
if (res != CUDA_SUCCESS) {
printf("cuMemMap (a) failed\n");
return -1;
}
res = cuMemGetPhysAddr(&a_phys, (void*)a_buf);
if (res != CUDA_SUCCESS) {
printf("cuMemGetPhysAddress (a) failed\n");
return -1;
}
/*printf("a[]: Physical Address 0x%llx\n", a_phys);*/
/* b[] */
res = cuMemAlloc(&b_dev, n*n * sizeof(unsigned int));
if (res != CUDA_SUCCESS) {
printf("cuMemAlloc (b) failed\n");
return -1;
}
res = cuMemMap((void**)&b_buf, b_dev, n*n * sizeof(unsigned int));
if (res != CUDA_SUCCESS) {
printf("cuMemMap (b) failed\n");
return -1;
}
res = cuMemGetPhysAddr(&b_phys, (void*)b_buf);
if (res != CUDA_SUCCESS) {
printf("cuMemGetPhysAddress (b) failed\n");
return -1;
}
/*printf("b[]: Physical Address 0x%llx\n", b_phys);*/
/* c[] */
res = cuMemAlloc(&c_dev, n*n * sizeof(unsigned int));
if (res != CUDA_SUCCESS) {
printf("cuMemAlloc (c) failed\n");
return -1;
}
res = cuMemMap((void**)&c_buf, c_dev, n*n * sizeof(unsigned int));
if (res != CUDA_SUCCESS) {
printf("cuMemMap (c) failed\n");
return -1;
}
res = cuMemGetPhysAddr(&c_phys, (void*)c_buf);
if (res != CUDA_SUCCESS) {
printf("cuMemGetPhysAddress (c) failed\n");
return -1;
}
/*printf("c[]: Physical Address 0x%llx\n", c_phys);*/
gettimeofday(&tv_data_init_start, NULL);
/* initialize A[] & B[] */
for (i = 0; i < n; i++) {
idx = i*n;
for(j = 0; j < n; j++) {
a_buf[idx++] = i;
}
}
for (i = 0; i < n; i++) {
idx = i*n;
for(j = 0; j < n; j++) {
b_buf[idx++] = i;
}
}
gettimeofday(&tv_h2d_start, NULL);
gettimeofday(&tv_h2d_end, NULL);
gettimeofday(&tv_conf_kern_start, NULL);
/* set kernel parameters */
res = cuParamSeti(function, 0, a_dev);
if (res != CUDA_SUCCESS) {
printf("cuParamSeti (a) failed: res = %lu\n", (unsigned long)res);
return -1;
}
res = cuParamSeti(function, 4, a_dev >> 32);
if (res != CUDA_SUCCESS) {
printf("cuParamSeti (a) failed: res = %lu\n", (unsigned long)res);
return -1;
}
res = cuParamSeti(function, 8, b_dev);
if (res != CUDA_SUCCESS) {
printf("cuParamSeti (b) failed: res = %lu\n", (unsigned long)res);
return -1;
}
res = cuParamSeti(function, 12, b_dev >> 32);
if (res != CUDA_SUCCESS) {
printf("cuParamSeti (b) failed: res = %lu\n", (unsigned long)res);
return -1;
}
res = cuParamSeti(function, 16, c_dev);
if (res != CUDA_SUCCESS) {
printf("cuParamSeti (c) failed: res = %lu\n", (unsigned long)res);
return -1;
}
res = cuParamSeti(function, 20, c_dev >> 32);
if (res != CUDA_SUCCESS) {
printf("cuParamSeti (c) failed: res = %lu\n", (unsigned long)res);
return -1;
}
res = cuParamSeti(function, 24, n);
if (res != CUDA_SUCCESS) {
printf("cuParamSeti (c) failed: res = %lu\n", (unsigned long)res);
return -1;
}
res = cuParamSetSize(function, 28);
if (res != CUDA_SUCCESS) {
printf("cuParamSetSize failed: res = %lu\n", (unsigned long)res);
return -1;
}
gettimeofday(&tv_exec_start, NULL);
/* launch the kernel */
res = cuLaunchGrid(function, grid_x, grid_y);
if (res != CUDA_SUCCESS) {
printf("cuLaunchGrid failed: res = %lu\n", (unsigned long)res);
return -1;
}
cuCtxSynchronize();
gettimeofday(&tv_exec_end, NULL);
gettimeofday(&tv_d2h_start, NULL);
/* download c[] */
memcpy(c, c_buf, n*n*sizeof(unsigned int));
gettimeofday(&tv_d2h_end, NULL);
/* Read back */
for (i = 0; i < n; i++) {
idx = i*n;
for(j = 0; j < n; j++) {
dummy_c = c[idx++];
}
}
gettimeofday(&tv_close_start, NULL);
res = cuMemUnmap((void*)a_buf);
if (res != CUDA_SUCCESS) {
printf("cuMemUnmap (a) failed: res = %lu\n", (unsigned long)res);
return -1;
}
res = cuMemFree(a_dev);
if (res != CUDA_SUCCESS) {
printf("cuMemFree (a) failed: res = %lu\n", (unsigned long)res);
return -1;
}
res = cuMemUnmap((void*)b_buf);
if (res != CUDA_SUCCESS) {
printf("cuMemUnmap (b) failed: res = %lu\n", (unsigned long)res);
return -1;
}
res = cuMemFree(b_dev);
if (res != CUDA_SUCCESS) {
printf("cuMemFree (b) failed: res = %lu\n", (unsigned long)res);
return -1;
}
res = cuMemUnmap((void*)c_buf);
if (res != CUDA_SUCCESS) {
printf("cuMemUnmap (c) failed: res = %lu\n", (unsigned long)res);
return -1;
}
res = cuMemFree(c_dev);
if (res != CUDA_SUCCESS) {
printf("cuMemFree (c) failed: res = %lu\n", (unsigned long)res);
return -1;
}
res = cuModuleUnload(module);
if (res != CUDA_SUCCESS) {
printf("cuModuleUnload failed: res = %lu\n", (unsigned long)res);
return -1;
}
res = cuCtxDestroy(ctx);
if (res != CUDA_SUCCESS) {
printf("cuCtxDestroy failed: res = %lu\n", (unsigned long)res);
return -1;
}
gettimeofday(&tv_total_end, NULL);
tvsub(&tv_mem_alloc_start, &tv_total_start, &tv);
init_gpu = tv.tv_sec * 1000.0 + (float) tv.tv_usec / 1000.0;
tvsub(&tv_data_init_start, &tv_mem_alloc_start, &tv);
mem_alloc = tv.tv_sec * 1000.0 + (float) tv.tv_usec / 1000.0;
tvsub(&tv_h2d_start, &tv_data_init_start, &tv);
data_init = tv.tv_sec * 1000.0 + (float) tv.tv_usec / 1000.0;
tvsub(&tv_h2d_end, &tv_h2d_start, &tv);
h2d = tv.tv_sec * 1000.0 + (float) tv.tv_usec / 1000.0;
tvsub(&tv_exec_start, &tv_conf_kern_start, &tv);
configure_kernel = tv.tv_sec * 1000.0 + (float) tv.tv_usec / 1000.0;
tvsub(&tv_exec_end, &tv_exec_start, &tv);
exec = tv.tv_sec * 1000.0 + (float) tv.tv_usec / 1000.0;
tvsub(&tv_d2h_end, &tv_d2h_start, &tv);
d2h = tv.tv_sec * 1000.0 + (float) tv.tv_usec / 1000.0;
tvsub(&tv_close_start, &tv_d2h_end, &tv);
data_read = tv.tv_sec * 1000.0 + (float) tv.tv_usec / 1000.0;
tvsub(&tv_total_end, &tv_close_start, &tv);
close_gpu = tv.tv_sec * 1000.0 + (float) tv.tv_usec / 1000.0;
tvsub(&tv_total_end, &tv_total_start, &tv);
total = tv.tv_sec * 1000.0 + (float) tv.tv_usec / 1000.0;
printf("Init: %f\n", init_gpu);
printf("MemAlloc: %f\n", mem_alloc);
printf("DataInit: %f\n", data_init);
printf("HtoD: %f\n", h2d);
printf("KernConf: %f\n", configure_kernel);
printf("Exec: %f\n", exec);
printf("DtoH: %f\n", d2h);
printf("DataRead: %f\n", data_read);
printf("Close: %f\n", close_gpu);
printf("Total: %f\n", total);
return ret;
}
JNIEXPORT void JNICALL Java_org_trifort_rootbeer_runtime_CUDAContext_cudaRun
(JNIEnv *env, jobject this_ref, jint device_index, jbyteArray cubin_file,
jint cubin_length, jint block_shape_x, jint grid_shape_x, jint num_threads,
jobject object_mem, jobject handles_mem, jobject exceptions_mem,
jobject class_mem)
{
CUresult status;
CUdevice device;
CUcontext context;
CUmodule module;
CUfunction function;
void * fatcubin;
int offset;
int info_space_size;
CUdeviceptr gpu_info_space;
CUdeviceptr gpu_object_mem;
CUdeviceptr gpu_handles_mem;
CUdeviceptr gpu_exceptions_mem;
CUdeviceptr gpu_class_mem;
CUdeviceptr gpu_heap_end;
CUdeviceptr gpu_buffer_size;
void * cpu_object_mem;
void * cpu_handles_mem;
void * cpu_exceptions_mem;
void * cpu_class_mem;
jlong cpu_object_mem_size;
jlong cpu_handles_mem_size;
jlong cpu_exceptions_mem_size;
jlong cpu_class_mem_size;
jlong cpu_heap_end;
jclass cuda_memory_class;
jmethodID get_address_method;
jmethodID get_size_method;
jmethodID get_heap_end_method;
jlong * info_space;
//----------------------------------------------------------------------------
//init device and function
//----------------------------------------------------------------------------
status = cuDeviceGet(&device, device_index);
CHECK_STATUS(env, "Error in cuDeviceGet", status, device)
status = cuCtxCreate(&context, CU_CTX_MAP_HOST, device);
CHECK_STATUS(env,"Error in cuCtxCreate", status, device)
fatcubin = malloc(cubin_length);
(*env)->GetByteArrayRegion(env, cubin_file, 0, cubin_length, fatcubin);
status = cuModuleLoadFatBinary(&module, fatcubin);
CHECK_STATUS(env, "Error in cuModuleLoad", status, device)
free(fatcubin);
status = cuModuleGetFunction(&function, module, "_Z5entryPcS_PiPxS1_S0_S0_i");
CHECK_STATUS(env, "Error in cuModuleGetFunction", status, device)
//----------------------------------------------------------------------------
//get handles from java
//----------------------------------------------------------------------------
cuda_memory_class = (*env)->FindClass(env, "org/trifort/rootbeer/runtime/FixedMemory");
get_address_method = (*env)->GetMethodID(env, cuda_memory_class, "getAddress", "()J");
get_size_method = (*env)->GetMethodID(env, cuda_memory_class, "getSize", "()J");
get_heap_end_method = (*env)->GetMethodID(env, cuda_memory_class, "getHeapEndPtr", "()J");
cpu_object_mem = (void *) (*env)->CallLongMethod(env, object_mem, get_address_method);
cpu_object_mem_size = (*env)->CallLongMethod(env, object_mem, get_size_method);
cpu_heap_end = (*env)->CallLongMethod(env, object_mem, get_heap_end_method);
cpu_handles_mem = (void *) (*env)->CallLongMethod(env, handles_mem, get_address_method);
cpu_handles_mem_size = (*env)->CallLongMethod(env, handles_mem, get_size_method);
cpu_exceptions_mem = (void *) (*env)->CallLongMethod(env, exceptions_mem, get_address_method);
cpu_exceptions_mem_size = (*env)->CallLongMethod(env, exceptions_mem, get_size_method);
cpu_class_mem = (void *) (*env)->CallLongMethod(env, class_mem, get_address_method);
cpu_class_mem_size = (*env)->CallLongMethod(env, class_mem, get_size_method);
info_space_size = 1024;
info_space = (jlong *) malloc(info_space_size);
info_space[1] = (*env)->CallLongMethod(env, object_mem, get_heap_end_method);
//----------------------------------------------------------------------------
//allocate mem
//----------------------------------------------------------------------------
status = cuMemAlloc(&gpu_info_space, info_space_size);
CHECK_STATUS(env, "Error in cuMemAlloc: gpu_info_mem", status, device)
status = cuMemAlloc(&gpu_object_mem, cpu_object_mem_size);
CHECK_STATUS(env, "Error in cuMemAlloc: gpu_object_mem", status, device)
status = cuMemAlloc(&gpu_handles_mem, cpu_handles_mem_size);
CHECK_STATUS(env, "Error in cuMemAlloc: gpu_handles_mem", status, device)
status = cuMemAlloc(&gpu_exceptions_mem, cpu_exceptions_mem_size);
CHECK_STATUS(env, "Error in cuMemAlloc: gpu_exceptions_mem", status, device)
status = cuMemAlloc(&gpu_class_mem, cpu_class_mem_size);
CHECK_STATUS(env, "Error in cuMemAlloc: gpu_class_mem", status, device)
status = cuMemAlloc(&gpu_heap_end, 8);
CHECK_STATUS(env, "Error in cuMemAlloc: gpu_heap_end", status, device)
status = cuMemAlloc(&gpu_buffer_size, 8);
CHECK_STATUS(env, "Error in cuMemAlloc: gpu_buffer_size", status, device)
//----------------------------------------------------------------------------
//set function parameters
//----------------------------------------------------------------------------
status = cuParamSetSize(function, (7 * sizeof(CUdeviceptr) + sizeof(int)));
CHECK_STATUS(env, "Error in cuParamSetSize", status, device)
offset = 0;
status = cuParamSetv(function, offset, (void *) &gpu_info_space, sizeof(CUdeviceptr));
CHECK_STATUS(env, "Error in cuParamSetv gpu_info_space", status, device)
offset += sizeof(CUdeviceptr);
status = cuParamSetv(function, offset, (void *) &gpu_object_mem, sizeof(CUdeviceptr));
CHECK_STATUS(env, "Error in cuParamSetv: gpu_object_mem", status, device)
offset += sizeof(CUdeviceptr);
status = cuParamSetv(function, offset, (void *) &gpu_handles_mem, sizeof(CUdeviceptr));
CHECK_STATUS(env, "Error in cuParamSetv: gpu_handles_mem %", status, device)
offset += sizeof(CUdeviceptr);
status = cuParamSetv(function, offset, (void *) &gpu_heap_end, sizeof(CUdeviceptr));
CHECK_STATUS(env, "Error in cuParamSetv: gpu_heap_end", status, device)
offset += sizeof(CUdeviceptr);
status = cuParamSetv(function, offset, (void *) &gpu_buffer_size, sizeof(CUdeviceptr));
CHECK_STATUS(env, "Error in cuParamSetv: gpu_buffer_size", status, device)
offset += sizeof(CUdeviceptr);
status = cuParamSetv(function, offset, (void *) &gpu_exceptions_mem, sizeof(CUdeviceptr));
CHECK_STATUS(env, "Error in cuParamSetv: gpu_exceptions_mem", status, device)
offset += sizeof(CUdeviceptr);
status = cuParamSetv(function, offset, (void *) &gpu_class_mem, sizeof(CUdeviceptr));
CHECK_STATUS(env, "Error in cuParamSetv: gpu_class_mem", status, device)
offset += sizeof(CUdeviceptr);
status = cuParamSeti(function, offset, num_threads);
CHECK_STATUS(env, "Error in cuParamSetv: num_threads", status, device)
offset += sizeof(int);
//----------------------------------------------------------------------------
//copy data
//----------------------------------------------------------------------------
status = cuMemcpyHtoD(gpu_info_space, info_space, info_space_size);
CHECK_STATUS(env, "Error in cuMemcpyHtoD: info_space", status, device)
status = cuMemcpyHtoD(gpu_object_mem, cpu_object_mem, cpu_object_mem_size);
CHECK_STATUS(env, "Error in cuMemcpyHtoD: gpu_object_mem", status, device)
status = cuMemcpyHtoD(gpu_handles_mem, cpu_handles_mem, cpu_handles_mem_size);
CHECK_STATUS(env, "Error in cuMemcpyHtoD: gpu_handles_mem", status, device)
status = cuMemcpyHtoD(gpu_class_mem, cpu_class_mem, cpu_class_mem_size);
CHECK_STATUS(env, "Error in cuMemcpyHtoD: gpu_class_mem", status, device)
status = cuMemcpyHtoD(gpu_heap_end, &cpu_heap_end, sizeof(jlong));
CHECK_STATUS(env, "Error in cuMemcpyHtoD: gpu_heap_end", status, device)
status = cuMemcpyHtoD(gpu_buffer_size, &cpu_object_mem_size, sizeof(jlong));
CHECK_STATUS(env, "Error in cuMemcpyHtoD: gpu_buffer_size", status, device)
status = cuMemcpyHtoD(gpu_exceptions_mem, cpu_exceptions_mem, cpu_exceptions_mem_size);
CHECK_STATUS(env, "Error in cuMemcpyDtoH: gpu_exceptions_mem", status, device)
//----------------------------------------------------------------------------
//launch
//----------------------------------------------------------------------------
status = cuFuncSetBlockShape(function, block_shape_x, 1, 1);
CHECK_STATUS(env, "Error in cuFuncSetBlockShape", status, device);
status = cuLaunchGrid(function, grid_shape_x, 1);
CHECK_STATUS(env, "Error in cuLaunchGrid", status, device)
status = cuCtxSynchronize();
CHECK_STATUS(env, "Error in cuCtxSynchronize", status, device)
//----------------------------------------------------------------------------
//copy data back
//----------------------------------------------------------------------------
status = cuMemcpyDtoH(info_space, gpu_info_space, info_space_size);
CHECK_STATUS(env, "Error in cuMemcpyDtoH: gpu_info_space", status, device)
cpu_heap_end = info_space[1];
status = cuMemcpyDtoH(cpu_object_mem, gpu_object_mem, cpu_heap_end);
CHECK_STATUS(env, "Error in cuMemcpyDtoH: gpu_object_mem", status, device)
status = cuMemcpyDtoH(cpu_exceptions_mem, gpu_exceptions_mem, cpu_exceptions_mem_size);
CHECK_STATUS(env, "Error in cuMemcpyDtoH: gpu_exceptions_mem", status, device)
//----------------------------------------------------------------------------
//free resources
//----------------------------------------------------------------------------
free(info_space);
cuMemFree(gpu_info_space);
cuMemFree(gpu_object_mem);
cuMemFree(gpu_handles_mem);
cuMemFree(gpu_exceptions_mem);
cuMemFree(gpu_class_mem);
cuMemFree(gpu_heap_end);
cuMemFree(gpu_buffer_size);
cuCtxDestroy(context);
}
void load_and_test(CUmodule cuModule, char * test_name)
{
try
{
CUfunction proc;
test(cuModuleGetFunction(&proc, cuModule, test_name), "cuModuleGetFunction");
int max = 1000;
bool * h_R = (bool*)malloc(max * sizeof(bool));
memset(h_R, 0, max * sizeof(bool));
CUdeviceptr d_R;
test(cuMemAlloc(&d_R, max * sizeof(bool)), "cuMemAlloc");
test(cuMemcpyHtoD(d_R, h_R, max * sizeof(bool)), "cuMemcpyHtoD");
CUdeviceptr d_N;
int h_N = 0;
test(cuMemAlloc(&d_N, sizeof(int)), "cuMemAlloc");
test(cuMemcpyHtoD(d_N, &h_N, sizeof(int)), "cuMemcpyHtoD");
int offset = 0;
void* ptr;
ptr = (void*)(size_t)d_R;
ALIGN_UP(offset, __alignof(ptr));
test(cuParamSetv(proc, offset, &ptr, sizeof(ptr)), "cuParamSetv");
offset += sizeof(ptr);
ptr = (void*)(size_t)d_N;
ALIGN_UP(offset, __alignof(ptr));
test(cuParamSetv(proc, offset, &ptr, sizeof(ptr)), "cuParamSetv");
offset += sizeof(ptr);
test(cuParamSetSize(proc, offset), "cuParamSetSize");
int threadsPerBlock = 1;
int blocksPerGrid = 1;
test(cuFuncSetBlockShape(proc, threadsPerBlock, 1, 1), "cuFuncSetBlockShape");
test(cuLaunchGrid(proc, blocksPerGrid, 1), "cuLaunchGrid");
test(cuMemcpyDtoH(h_R, d_R, max * sizeof(bool)), "cuMemcpyDtoH");
test(cuMemcpyDtoH(&h_N, d_N, sizeof(int)), "cuMemcpyDtoH");
test(cuMemFree(d_R), "cuMemFree");
test(cuMemFree(d_N), "cuMemFree");
bool failed = false;
for (int i = 0; i < h_N; ++i)
{
if (h_R[i] == 0)
{
failed = true;
std::cout << "\nTest " << i << " failed.\n";
std::cout.flush();
}
}
if (! failed)
std::cout << test_name << " passed.\n";
else {
std::cout << test_name << " failed.\n";
}
}
catch (...)
{
std::string s = test_name;
s = s.append(" crashed.\n");
test(1, s.c_str());
}
}
int cuda_test_fmadd(unsigned int n, char *path)
{
int i, j, idx;
CUresult res;
CUdevice dev;
CUcontext ctx;
CUfunction function;
CUmodule module;
CUdeviceptr a_dev, b_dev, c_dev;
float *a = (float *) malloc (n*n * sizeof(float));
float *b = (float *) malloc (n*n * sizeof(float));
float *c = (float *) malloc (n*n * sizeof(float));
int block_x, block_y, grid_x, grid_y;
int offset;
char fname[256];
struct timeval tv;
struct timeval tv_total_start, tv_total_end;
float total;
struct timeval tv_h2d_start, tv_h2d_end;
float h2d;
struct timeval tv_d2h_start, tv_d2h_end;
float d2h;
struct timeval tv_exec_start, tv_exec_end;
float exec;
/* initialize A[] & B[] */
for (i = 0; i < n; i++) {
for(j = 0; j < n; j++) {
idx = i * n + j;
a[idx] = i + 0.1;
b[idx] = i + 0.1;
}
}
/* block_x * block_y should not exceed 512. */
block_x = n < 16 ? n : 16;
block_y = n < 16 ? n : 16;
grid_x = n / block_x;
if (n % block_x != 0)
grid_x++;
grid_y = n / block_y;
if (n % block_y != 0)
grid_y++;
printf("block = (%d, %d)\n", block_x, block_y);
printf("grid = (%d, %d)\n", grid_x, grid_y);
gettimeofday(&tv_total_start, NULL);
res = cuInit(0);
if (res != CUDA_SUCCESS) {
printf("cuInit failed: res = %lu\n", (unsigned long)res);
return -1;
}
res = cuDeviceGet(&dev, 0);
if (res != CUDA_SUCCESS) {
printf("cuDeviceGet failed: res = %lu\n", (unsigned long)res);
return -1;
}
res = cuCtxCreate(&ctx, 0, dev);
if (res != CUDA_SUCCESS) {
printf("cuCtxCreate failed: res = %lu\n", (unsigned long)res);
return -1;
}
sprintf(fname, "%s/fmadd_gpu.cubin", path);
res = cuModuleLoad(&module, fname);
if (res != CUDA_SUCCESS) {
printf("cuModuleLoad() failed\n");
return -1;
}
res = cuModuleGetFunction(&function, module, "_Z3addPfS_S_i");
if (res != CUDA_SUCCESS) {
printf("cuModuleGetFunction() failed\n");
return -1;
}
res = cuFuncSetSharedSize(function, 0x40); /* just random */
if (res != CUDA_SUCCESS) {
printf("cuFuncSetSharedSize() failed\n");
return -1;
}
res = cuFuncSetBlockShape(function, block_x, block_y, 1);
if (res != CUDA_SUCCESS) {
printf("cuFuncSetBlockShape() failed\n");
return -1;
}
/* a[] */
res = cuMemAlloc(&a_dev, n*n * sizeof(float));
if (res != CUDA_SUCCESS) {
printf("cuMemAlloc (a) failed\n");
return -1;
}
/* b[] */
res = cuMemAlloc(&b_dev, n*n * sizeof(float));
if (res != CUDA_SUCCESS) {
printf("cuMemAlloc (b) failed\n");
return -1;
}
/* c[] */
res = cuMemAlloc(&c_dev, n*n * sizeof(float));
if (res != CUDA_SUCCESS) {
printf("cuMemAlloc (c) failed\n");
return -1;
}
gettimeofday(&tv_h2d_start, NULL);
/* upload a[] and b[] */
res = cuMemcpyHtoD(a_dev, a, n*n * sizeof(float));
if (res != CUDA_SUCCESS) {
printf("cuMemcpyHtoD (a) failed: res = %lu\n", (unsigned long)res);
return -1;
}
res = cuMemcpyHtoD(b_dev, b, n*n * sizeof(float));
if (res != CUDA_SUCCESS) {
printf("cuMemcpyHtoD (b) failed: res = %lu\n", (unsigned long)res);
return -1;
}
gettimeofday(&tv_h2d_end, NULL);
/* set kernel parameters */
offset = 0;
res = cuParamSetv(function, offset, &a_dev, sizeof(a_dev));
if (res != CUDA_SUCCESS) {
printf("cuParamSeti (a) failed: res = %lu\n", (unsigned long)res);
return -1;
}
offset += sizeof(a_dev);
res = cuParamSetv(function, offset, &b_dev, sizeof(b_dev));
if (res != CUDA_SUCCESS) {
printf("cuParamSeti (b) failed: res = %lu\n", (unsigned long)res);
return -1;
}
offset += sizeof(b_dev);
res = cuParamSetv(function, offset, &c_dev, sizeof(c_dev));
if (res != CUDA_SUCCESS) {
printf("cuParamSeti (c) failed: res = %lu\n", (unsigned long)res);
return -1;
}
offset += sizeof(c_dev);
res = cuParamSetv(function, offset, &n, sizeof(n));
if (res != CUDA_SUCCESS) {
printf("cuParamSeti (c) failed: res = %lu\n", (unsigned long)res);
return -1;
}
offset += sizeof(n);
res = cuParamSetSize(function, offset);
if (res != CUDA_SUCCESS) {
printf("cuParamSetSize failed: res = %lu\n", (unsigned long)res);
return -1;
}
gettimeofday(&tv_exec_start, NULL);
/* launch the kernel */
res = cuLaunchGrid(function, grid_x, grid_y);
if (res != CUDA_SUCCESS) {
printf("cuLaunchGrid failed: res = %lu\n", (unsigned long)res);
return -1;
}
cuCtxSynchronize();
gettimeofday(&tv_exec_end, NULL);
gettimeofday(&tv_d2h_start, NULL);
/* download c[] */
res = cuMemcpyDtoH(c, c_dev, n*n * sizeof(float));
if (res != CUDA_SUCCESS) {
printf("cuMemcpyDtoH (c) failed: res = %lu\n", (unsigned long)res);
return -1;
}
gettimeofday(&tv_d2h_end, NULL);
res = cuMemFree(a_dev);
if (res != CUDA_SUCCESS) {
printf("cuMemFree (a) failed: res = %lu\n", (unsigned long)res);
return -1;
}
res = cuMemFree(b_dev);
if (res != CUDA_SUCCESS) {
printf("cuMemFree (b) failed: res = %lu\n", (unsigned long)res);
return -1;
}
res = cuMemFree(c_dev);
if (res != CUDA_SUCCESS) {
printf("cuMemFree (c) failed: res = %lu\n", (unsigned long)res);
return -1;
}
res = cuModuleUnload(module);
if (res != CUDA_SUCCESS) {
printf("cuModuleUnload failed: res = %lu\n", (unsigned long)res);
return -1;
}
res = cuCtxDestroy(ctx);
if (res != CUDA_SUCCESS) {
printf("cuCtxDestroy failed: res = %lu\n", (unsigned long)res);
return -1;
}
gettimeofday(&tv_total_end, NULL);
/* check the results */
i = j = idx = 0;
while (i < n) {
while (j < n) {
idx = i * n + j;
if (c[idx] != a[idx] + b[idx]) {
printf("c[%d] = %f\n", idx, c[idx]);
printf("a[%d]+b[%d] = %f\n", idx, idx, a[idx]+b[idx]);
return -1;
}
j++;
}
i++;
}
free(a);
free(b);
free(c);
tvsub(&tv_h2d_end, &tv_h2d_start, &tv);
h2d = tv.tv_sec * 1000.0 + (float) tv.tv_usec / 1000.0;
tvsub(&tv_d2h_end, &tv_d2h_start, &tv);
d2h = tv.tv_sec * 1000.0 + (float) tv.tv_usec / 1000.0;
tvsub(&tv_exec_end, &tv_exec_start, &tv);
exec = tv.tv_sec * 1000.0 + (float) tv.tv_usec / 1000.0;
tvsub(&tv_total_end, &tv_total_start, &tv);
total = tv.tv_sec * 1000.0 + (float) tv.tv_usec / 1000.0;
printf("HtoD: %f\n", h2d);
printf("DtoH: %f\n", d2h);
printf("Exec: %f\n", exec);
printf("Time (Memcpy + Launch): %f\n", h2d + d2h + exec);
printf("Total: %f\n", total);
return 0;
}
/**
* Invokes the kernel f on a 1 x 1 x 1 grid of blocks. The block contains the
* number of threads specified by a previous call to cuFuncSetBlockShape().
*
* Parameters:
* f - Kernel to launch
*
* Returns:
* CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
* CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE,
* CUDA_ERROR_LAUNCH_FAILED, CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES,
* CUDA_ERROR_LAUNCH_TIMEOUT, CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING
*/
CUresult cuLaunch(CUfunction f)
{
return cuLaunchGrid(f, 1, 1);
}
int gib_recover ( void *buffers, int buf_size, int *buf_ids, int recover_last,
gib_context c ) {
ERROR_CHECK_FAIL(cuCtxPushCurrent(((gpu_context)(c->acc_context))->pCtx));
#if !GIB_USE_MMAP
if (buf_size > gib_buf_size) {
int rc = gib_cpu_recover(buffers, buf_size, buf_ids, recover_last, c);
ERROR_CHECK_FAIL(cuCtxPopCurrent(&((gpu_context)(c->acc_context))->pCtx));
return rc;
}
#endif
int i, j;
int n = c->n;
int m = c->m;
unsigned char A[128*128], inv[128*128], modA[128*128];
for (i = n; i < n+recover_last; i++)
if (buf_ids[i] >= n) {
fprintf(stderr, "Attempting to recover a parity buffer, not allowed\n");
return GIB_ERR;
}
gib_galois_gen_A(A, m+n, n);
/* Modify the matrix to have the failed drives reflected */
for (i = 0; i < n; i++)
for (j = 0; j < n; j++)
modA[i*n+j] = A[buf_ids[i]*n+j];
gib_galois_gaussian_elim(modA, inv, n, n);
/* Copy row buf_ids[i] into row i */
for (i = n; i < n+recover_last; i++)
for (j = 0; j < n; j++)
modA[i*n+j] = inv[buf_ids[i]*n+j];
int nthreads_per_block = 128;
int fetch_size = sizeof(int)*nthreads_per_block;
int nblocks = (buf_size + fetch_size - 1)/fetch_size;
gpu_context gpu_c = (gpu_context) c->acc_context;
CUdeviceptr F_d;
ERROR_CHECK_FAIL(cuModuleGetGlobal(&F_d, NULL, gpu_c->module, "F_d"));
ERROR_CHECK_FAIL(cuMemcpyHtoD(F_d, modA+n*n, (c->m)*(c->n)));
#if !GIB_USE_MMAP
ERROR_CHECK_FAIL(cuMemcpyHtoD(gpu_c->buffers, buffers, (c->n)*buf_size));
#endif
ERROR_CHECK_FAIL(cuFuncSetBlockShape(gpu_c->recover, nthreads_per_block,
1, 1));
int offset = 0;
void *ptr;
#if GIB_USE_MMAP
CUdeviceptr cpu_buffers;
ERROR_CHECK_FAIL(cuMemHostGetDevicePointer(&cpu_buffers, buffers, 0));
ptr = (void *)cpu_buffers;
#else
ptr = (void *)gpu_c->buffers;
#endif
ERROR_CHECK_FAIL(cuParamSetv(gpu_c->recover, offset, &ptr, sizeof(ptr)));
offset += sizeof(ptr);
ERROR_CHECK_FAIL(cuParamSetv(gpu_c->recover, offset, &buf_size,
sizeof(buf_size)));
offset += sizeof(buf_size);
ERROR_CHECK_FAIL(cuParamSetv(gpu_c->recover, offset, &recover_last,
sizeof(recover_last)));
offset += sizeof(recover_last);
ERROR_CHECK_FAIL(cuParamSetSize(gpu_c->recover, offset));
ERROR_CHECK_FAIL(cuLaunchGrid(gpu_c->recover, nblocks, 1));
#if !GIB_USE_MMAP
CUdeviceptr tmp_d = gpu_c->buffers + c->n*buf_size;
void *tmp_h = (void *)((unsigned char *)(buffers) + c->n*buf_size);
ERROR_CHECK_FAIL(cuMemcpyDtoH(tmp_h, tmp_d, recover_last*buf_size));
#else
cuCtxSynchronize();
#endif
ERROR_CHECK_FAIL(cuCtxPopCurrent(&((gpu_context)(c->acc_context))->pCtx));
return GIB_SUC;
}
int gib_generate ( void *buffers, int buf_size, gib_context c ) {
ERROR_CHECK_FAIL(cuCtxPushCurrent(((gpu_context)(c->acc_context))->pCtx));
/* Do it all at once if the buffers are small enough */
#if !GIB_USE_MMAP
/* This is too large to do at once in the GPU memory we have allocated.
* Split it into several noncontiguous jobs.
*/
if (buf_size > gib_buf_size) {
int rc = gib_generate_nc(buffers, buf_size, buf_size, c);
ERROR_CHECK_FAIL(cuCtxPopCurrent(&((gpu_context)(c->acc_context))->pCtx));
return rc;
}
#endif
int nthreads_per_block = 128;
int fetch_size = sizeof(int)*nthreads_per_block;
int nblocks = (buf_size + fetch_size - 1)/fetch_size;
gpu_context gpu_c = (gpu_context) c->acc_context;
unsigned char F[256*256];
gib_galois_gen_F(F, c->m, c->n);
CUdeviceptr F_d;
ERROR_CHECK_FAIL(cuModuleGetGlobal(&F_d, NULL, gpu_c->module, "F_d"));
ERROR_CHECK_FAIL(cuMemcpyHtoD(F_d, F, (c->m)*(c->n)));
#if !GIB_USE_MMAP
/* Copy the buffers to memory */
ERROR_CHECK_FAIL(cuMemcpyHtoD(gpu_c->buffers, buffers,
(c->n)*buf_size));
#endif
/* Configure and launch */
ERROR_CHECK_FAIL(cuFuncSetBlockShape(gpu_c->checksum, nthreads_per_block,
1, 1));
int offset = 0;
void *ptr;
#if GIB_USE_MMAP
CUdeviceptr cpu_buffers;
ERROR_CHECK_FAIL(cuMemHostGetDevicePointer(&cpu_buffers, buffers, 0));
ptr = (void *)cpu_buffers;
#else
ptr = (void *)(gpu_c->buffers);
#endif
ERROR_CHECK_FAIL(cuParamSetv(gpu_c->checksum, offset, &ptr, sizeof(ptr)));
offset += sizeof(ptr);
ERROR_CHECK_FAIL(cuParamSetv(gpu_c->checksum, offset, &buf_size,
sizeof(buf_size)));
offset += sizeof(buf_size);
ERROR_CHECK_FAIL(cuParamSetSize(gpu_c->checksum, offset));
ERROR_CHECK_FAIL(cuLaunchGrid(gpu_c->checksum, nblocks, 1));
/* Get the results back */
#if !GIB_USE_MMAP
CUdeviceptr tmp_d = gpu_c->buffers + c->n*buf_size;
void *tmp_h = (void *)((unsigned char *)(buffers) + c->n*buf_size);
ERROR_CHECK_FAIL(cuMemcpyDtoH(tmp_h, tmp_d, (c->m)*buf_size));
#else
ERROR_CHECK_FAIL(cuCtxSynchronize());
#endif
ERROR_CHECK_FAIL(cuCtxPopCurrent(&((gpu_context)(c->acc_context))->pCtx));
return GIB_SUC;
}
// Host code
int main()
{
int N = 3;
size_t size = N * sizeof(float);
float* h_A = (float*)malloc(size);
float* h_B = (float*)malloc(size);
float* h_C = (float*)malloc(size);
// Set up vectors.
for (int i = 0; i < N; ++i)
{
h_A[i] = i * 1.0;
h_B[i] = i * 1.0 + 1;
h_C[i] = 0;
printf("i %d A %f B %f C %f\n", i, h_A[i], h_B[i], h_C[i]);
}
// Initialize
if (cuInit(0) != CUDA_SUCCESS)
exit (0);
// Get number of devices supporting CUDA
int deviceCount = 0;
cuDeviceGetCount(&deviceCount);
if (deviceCount == 0)
{
printf("There is no device supporting CUDA.\n");
exit (0);
}
// Get handle for device 0
CUdevice cuDevice = 0;
CUresult r1 = cuDeviceGet(&cuDevice, 0);
// Create context
CUcontext cuContext;
cuCtxCreate(&cuContext, 0, cuDevice);
// Create module from binary file
CUmodule cuModule;
CUresult r2 = cuModuleLoad(&cuModule, "VecAdd.ptx");
// Get function handle from module
CUfunction vecAdd;
CUresult r3 = cuModuleGetFunction(&vecAdd, cuModule, "VecAdd");
// Allocate vectors in device memory
CUdeviceptr d_A;
CUresult r4 = cuMemAlloc(&d_A, size);
CUdeviceptr d_B;
CUresult r5 = cuMemAlloc(&d_B, size);
CUdeviceptr d_C;
CUresult r6 = cuMemAlloc(&d_C, size);
// Copy vectors from host memory to device memory
// h_A and h_B are input vectors stored in host memory
CUresult r7 = cuMemcpyHtoD(d_A, h_A, size);
CUresult r8 = cuMemcpyHtoD(d_B, h_B, size);
// Invoke kernel
#define ALIGN_UP(offset, alignment) (offset) = ((offset) + (alignment) - 1) & ~((alignment) - 1)
int offset = 0;
void* ptr;
ptr = (void*)(size_t)d_A;
ALIGN_UP(offset, __alignof(ptr));
CUresult r9 = cuParamSetv(vecAdd, offset, &ptr, sizeof(ptr));
offset += sizeof(ptr);
ptr = (void*)(size_t)d_B;
ALIGN_UP(offset, __alignof(ptr));
CUresult r10 = cuParamSetv(vecAdd, offset, &ptr, sizeof(ptr));
offset += sizeof(ptr);
ptr = (void*)(size_t)d_C;
ALIGN_UP(offset, __alignof(ptr));
CUresult r11 = cuParamSetv(vecAdd, offset, &ptr, sizeof(ptr));
offset += sizeof(ptr);
ptr = (void*)(int)N;
ALIGN_UP(offset, __alignof(ptr));
CUresult r11a = cuParamSetv(vecAdd, offset, &ptr, sizeof(ptr));
offset += sizeof(ptr);
CUresult r12 = cuParamSetSize(vecAdd, offset);
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
CUresult r13 = cuFuncSetBlockShape(vecAdd, threadsPerBlock, 1, 1);
CUresult r14 = cuLaunchGrid(vecAdd, blocksPerGrid, 1);
// Copy result from device memory to host memory
// h_C contains the result in host memory
CUresult r15 = cuMemcpyDtoH(h_C, d_C, size);
for (int i = 0; i < N; ++i)
{
printf("i %d A %f B %f C %f\n", i, h_A[i], h_B[i], h_C[i]);
}
// Free device memory
cuMemFree(d_A);
cuMemFree(d_B);
cuMemFree(d_C);
}
void CUDARunner::FindBestConfiguration()
{
unsigned long lowb=16;
unsigned long highb=128;
unsigned long lowt=16;
unsigned long hight=256;
unsigned long bestb=16;
unsigned long bestt=16;
int offset=0;
void *ptr=0;
int64 besttime=std::numeric_limits<int64>::max();
if(m_requestedgrid>0 && m_requestedgrid<=65536)
{
lowb=m_requestedgrid;
highb=m_requestedgrid;
}
if(m_requestedthreads>0 && m_requestedthreads<=65536)
{
lowt=m_requestedthreads;
hight=m_requestedthreads;
}
for(int numb=lowb; numb<=highb; numb*=2)
{
for(int numt=lowt; numt<=hight; numt*=2)
{
if(AllocateResources(numb,numt)==true)
{
// clear out any existing error
CUresult err=CUDA_SUCCESS;
int64 st=GetTimeMillis();
for(int it=0; it<128*256*2 && err==CUDA_SUCCESS; it+=(numb*numt))
{
cuMemcpyHtoD(m_devin,m_in,sizeof(cuda_in));
offset=0;
int loops=64;
int bits=5;
ptr=(void *)(size_t)m_devin;
ALIGN_UP(offset, __alignof(ptr));
cuParamSetv(m_function,offset,&ptr,sizeof(ptr));
offset+=sizeof(ptr);
ptr=(void *)(size_t)m_devout;
ALIGN_UP(offset, __alignof(ptr));
cuParamSetv(m_function,offset,&ptr,sizeof(ptr));
offset+=sizeof(ptr);
ALIGN_UP(offset, __alignof(loops));
cuParamSeti(m_function,offset,loops);
offset+=sizeof(loops);
ALIGN_UP(offset, __alignof(bits));
cuParamSeti(m_function,offset,bits);
offset+=sizeof(bits);
cuParamSetSize(m_function,offset);
err=cuFuncSetBlockShape(m_function,numt,1,1);
if(err!=CUDA_SUCCESS)
{
printf("cuFuncSetBlockShape error %d\n",err);
continue;
}
err=cuLaunchGrid(m_function,numb,1);
if(err!=CUDA_SUCCESS)
{
printf("cuLaunchGrid error %d\n",err);
continue;
}
cuMemcpyDtoH(m_out,m_devout,numt*numb*sizeof(cuda_out));
if(err!=CUDA_SUCCESS)
{
printf("CUDA error %d\n",err);
}
}
int64 et=GetTimeMillis();
printf("Finding best configuration step end (%d,%d) %"PRI64d"ms prev best=%"PRI64d"ms\n",numb,numt,et-st,besttime);
if((et-st)<besttime && err==CUDA_SUCCESS)
{
bestb=numb;
bestt=numt;
besttime=et-st;
}
}
}
}
m_numb=bestb;
m_numt=bestt;
AllocateResources(m_numb,m_numt);
}
void Function::launch(int gridWidth, int gridHeight) const
{
detail::error_check(cuLaunchGrid(impl->func, gridWidth, gridHeight),
"Can't launch Cuda function grid");
}
int main() {
const int N = 9;
int bytes = sizeof(int) * N;
int *A_cpu = 0;
CUdeviceptr A_gpu = 0;
A_cpu = new int[N];
for (int i = 0; i < N; i++) {
A_cpu[i] = -1;
}
CUresult result = cuInit(0);
if (result != CUDA_SUCCESS) {
report("cuInit() failed: " << result);
return 1;
}
int driverVersion = 0;
result = cuDriverGetVersion(&driverVersion);
if (result != CUDA_SUCCESS) {
report("cuDriverGetVersion() failed: " << result);
}
int count = 0;
result = cuDeviceGetCount(&count);
if (result != CUDA_SUCCESS) {
report("cuDeviceGetCount() failed: " << result);
return 1;
}
CUdevice device;
result = cuDeviceGet(&device, 0);
if (result != CUDA_SUCCESS) {
report("cuDeviceGet() failed: " << result);
return 1;
}
char devName[256] = {0};
result = cuDeviceGetName(devName, 255, device);
if (result != CUDA_SUCCESS) {
report("cuDeviceGetName() failed: " << result);
return 1;
}
int major, minor;
result = cuDeviceComputeCapability(&major, &minor, device);
if (result != CUDA_SUCCESS) {
report("cuDeviceComputeCapability() failed: " << result);
return 1;
}
CUcontext ctx;
CUmodule module;
CUfunction function;
result = cuCtxCreate(&ctx, 0, device);
if (result != CUDA_SUCCESS) {
report("cuCtxCreate() failed: " << result);
return 1;
}
int pi = 0;
result = cuDeviceGetAttribute(&pi, CU_DEVICE_ATTRIBUTE_COMPUTE_MODE, device);
if (result != CUDA_SUCCESS) {
report("cuDeviceGetAttribute() failed: " << result);
}
result = cuModuleLoad(&module, "ocelot/cuda/test/driver/generic.ptx");
if (result != CUDA_SUCCESS) {
report("cuModuleLoad() failed: " << result);
return 1;
}
result = cuModuleGetFunction(&function, module, "genericmemory");
if (result != CUDA_SUCCESS) {
report("cuModuleGetFunction() failed: " << result);
return 1;
}
result = cuMemAlloc(&A_gpu, bytes);
if (result != CUDA_SUCCESS) {
report("cuMemAlloc() failed: " << result);
return 1;
}
result = cuMemcpyHtoD(A_gpu, A_cpu, bytes);
if (result != CUDA_SUCCESS) {
report("cuMemcpyHtoD() failed: " << result);
return 1;
}
struct {
int *A;
} parameters;
result = cuParamSetSize(function, sizeof(parameters));
if (result != CUDA_SUCCESS) {
report("cuParamSetSize() failed: " << result);
return 1;
}
parameters.A = reinterpret_cast<int *>(A_gpu);
result = cuParamSetv(function, 0, ¶meters.A, sizeof(parameters.A));
if (result != CUDA_SUCCESS) {
report("cuParamSetv() failed: " << result);
return 1;
}
result = cuFuncSetBlockShape(function, 1, 1, 1);
if (result != CUDA_SUCCESS) {
report("cuFuncSetBlockShape() failed: " << result);
return 1;
}
result = cuLaunchGrid(function, 1, 1);
if (result != CUDA_SUCCESS) {
report("cuLaunchGrid() failed: " << result);
return 1;
}
result = cuMemcpyDtoH(A_cpu, A_gpu, bytes);
if (result != CUDA_SUCCESS) {
report("cuMemcpyDtoH() failed: " << result);
return 1;
}
cuModuleUnload(module);
cuCtxDestroy(ctx);
int errors = 0;
for (int i = 0; i < 9; i++) {
if (i < 3 && !A_cpu[i] || i >= 3 && A_cpu[i]) {
++errors;
std::cout << "%p" << i << " - " << A_cpu[i] << "\n";
}
}
delete [] A_cpu;
std::cout << "Pass/Fail : " << (!errors ? "Pass" : "Fail") << std::endl;
return 0;
}
////////////////////////////////////////////////////////////////////////////////
//! Run a simple test for CUDA
////////////////////////////////////////////////////////////////////////////////
void
runTest(int argc, char** argv)
{
CUcontext cuContext;
// initialize CUDA
CUfunction pk = NULL;
const char cubin_name [] = "pass_kernel.cubin";
const char kernel_name [] = "pass_kernel";
CU_SAFE_CALL(initCuda(cuContext, argv[0], &pk, argc, argv, cubin_name, kernel_name));
printf("initCuda-returned CUfunction:\n");
// cuParamSetx, x=i f v
// http://visionexperts.blogspot.com/2010/07/cuda-parameter-alignment.html - check alignment
#define ALIGN_UP(offset, alignment) \
(offset) = ((offset) + (alignment) - 1) & ~((alignment) - 1)
size_t offset = 0;
// input integers
// CU paramset i.
for(int i = 0 ; i < NUM_ARG ; i++)
{
int align = __alignof(int);
ALIGN_UP(offset, align);
cuParamSeti(pk, offset, i);
printf ("offset %d = %d\n", i, offset);
offset += sizeof(int);
}
// return array for updated inputs
int size_int = sizeof(int);
int size_array = size_int * NUM_ARG;
CUdeviceptr d_return_values;
cuMemAlloc (&d_return_values, size_array);
void* ptr = (void*)(size_t)d_return_values;
int align = __alignof(ptr);
ALIGN_UP(offset, align);
cuParamSetv(pk, offset, &ptr, sizeof(ptr));
printf("return values offset:%d\n", offset);
offset += sizeof(ptr);
CUdeviceptr d_return_N;
cuMemAlloc(&d_return_N, size_int);
void* ptrN = (void*)(size_t)d_return_N;
int alignN = __alignof(ptrN);
ALIGN_UP(offset, alignN);
cuParamSetv(pk, offset, &ptrN, sizeof(ptr));
printf("return int offset:%d\n", offset);
offset += sizeof(ptrN);
// Calling kernel
int BLOCK_SIZE_X = NUM_ARG;
int BLOCK_SIZE_Y = 1;
int BLOCK_SIZE_Z = 1;
int GRID_SIZE = 1;
cutilDrvSafeCallNoSync(cuFuncSetBlockShape(pk, BLOCK_SIZE_X, BLOCK_SIZE_Y, BLOCK_SIZE_Z));
printf("paramsetsize:%d\n", offset);
CU_SAFE_CALL(cuParamSetSize(pk, offset));
CU_SAFE_CALL(cuLaunchGrid(pk, GRID_SIZE, GRID_SIZE));
int* h_return_values = (int*)malloc(NUM_ARG * sizeof(int));
CU_SAFE_CALL(cuMemcpyDtoH((void*)h_return_values, d_return_values, size_array));
CU_SAFE_CALL(cuMemFree(d_return_values));
for(int i=0;i<NUM_ARG;i++)
printf("%dth value = %d\n", i, h_return_values[i]);
free(h_return_values);
int* h_return_N = (int*)malloc(sizeof(int));
CU_SAFE_CALL(cuMemcpyDtoH((void*)h_return_N, d_return_N, size_int));
CU_SAFE_CALL(cuMemFree(d_return_N));
printf("%d sizeof array\n", *h_return_N);
if(cuContext !=NULL) cuCtxDetach(cuContext);
}
int main(int argc, char *argv[])
{
srand(time(NULL));
for(int k=0;k<4;k++)
{
int n = 30*(k+1);
float x = ((float) rand()) / (float) RAND_MAX;
float *a = new float[n+1];
float resultGPU;
for(int i = 0; i < n + 1; i++)
a[i] = i * 0.5*((float) rand()) / (float) RAND_MAX;
int blocks = (n + 1) / BLK_SZ;
if((n + 1) % BLK_SZ)
blocks++;
CUdevice hDevice;
CUcontext hContext;
CUmodule hModule;
CUfunction hFunction;
CALL( cuInit(0) );
CALL( cuDeviceGet(&hDevice, 0) );
CALL( cuCtxCreate(&hContext, 0, hDevice) );
CALL( cuModuleLoad(&hModule, "kernel.cubin") );
CALL( cuModuleGetFunction(&hFunction, hModule, "Polynomial") );
//dane wejsciowe - kopiowanie
CUdeviceptr DevA, DevResult;
CALL( cuMemAlloc(&DevA, (n+1)*sizeof(float) ) );
CALL( cuMemAlloc(&DevResult, sizeof(float) ) );
CALL( cuMemcpyHtoD(DevA, a, (n+1)*sizeof(float) ) );
CALL( cuFuncSetBlockShape(hFunction, BLK_SZ, 1, 1) );
//przekazanie parametrow do kernela
int offset = 0;
void *ptr;
ptr = (void*)(size_t)DevResult;
ALIGN_UP(offset, __alignof(ptr));
CALL( cuParamSetv(hFunction, offset, &ptr, sizeof(ptr)) );
offset += sizeof(ptr);
ptr = (void*)(size_t)DevA;
ALIGN_UP(offset, __alignof(ptr));
CALL( cuParamSetv(hFunction, offset, &ptr, sizeof(ptr)) );
offset += sizeof(ptr);
ALIGN_UP(offset, __alignof(float));
CALL( cuParamSetf(hFunction, offset, x) );
offset += sizeof(float);
ALIGN_UP(offset, __alignof(int));
CALL( cuParamSeti(hFunction, offset, n) );
offset += sizeof(int);
CALL( cuParamSetSize(hFunction, offset) );
CALL( cuLaunchGrid(hFunction, blocks, 1) );
//kopiowanie wyniku na hosta
CALL( cuMemcpyDtoH((void *) &resultGPU, DevResult, sizeof(float) ) );
//zwalnianie pamieci na urzadzeniu
CALL( cuMemFree(DevA) );
CALL( cuMemFree(DevResult) );
//obliczenia na CPU
float resultCPU = PolynomialCPU(a, x, n);
std::cout << "GPU:\t" << resultGPU << std::endl;
std::cout << "CPU:\t" << resultCPU << std::endl;
std::cout << "roznica:\t" << fabs(resultGPU - resultCPU) << std::endl;
delete [] a;
}
return 0;
}
int main(void)
{
// Initialize
if (cuInit(0) != CUDA_SUCCESS)
exit (0);
// Get number of devices supporting CUDA
int deviceCount = 0;
cuDeviceGetCount(&deviceCount);
if (deviceCount == 0) {
printf("There is no device supporting CUDA.\n");
exit (0);
}
// Get handle for device 0
CUdevice cuDevice = 0;
cuDeviceGet(&cuDevice, 0);
// Create context
CUcontext cuContext;
cuCtxCreate(&cuContext, 0, cuDevice);
// Create module from binary file
CUmodule cuModule;
cuModuleLoad(&cuModule, “VecAdd.ptx”);
// Get function handle from module
CUfunction vecAdd;
cuModuleGetFunction(&vecAdd, cuModule, "VecAdd");
// Allocate vectors in device memory
size_t size = N * sizeof(float);
CUdeviceptr d_A;
cuMemAlloc(&d_A, size);
CUdeviceptr d_B;
cuMemAlloc(&d_B, size);
CUdeviceptr d_C;
cuMemAlloc(&d_C, size);
// Copy vectors from host memory to device memory
// h_A and h_B are input vectors stored in host memory
cuMemcpyHtoD(d_A, h_A, size);
cuMemcpyHtoD(d_B, h_B, size);
// Invoke kernel
#define ALIGN_UP(offset, alignment) \
(offset) = ((offset) + (alignment) – 1) & ~((alignment) – 1)
int offset = 0;
ALIGN_UP(offset, __alignof(d_A));
cuParamSetv(vecAdd, offset, &d_A, sizeof(d_A));
offset += sizeof(d_A);
ALIGN_UP(offset, __alignof(d_B));
cuParamSetv(vecAdd, offset, &d_B, sizeof(d_B));
offset += sizeof(d_B);
ALIGN_UP(offset, __alignof(d_C));
cuParamSetv(vecAdd, offset, &d_C, sizeof(d_C));
offset += sizeof(d_C);
cuParamSetSize(VecAdd, offset);
int threadsPerBlock = 256;
int blocksPerGrid =
(N + threadsPerBlock – 1) / threadsPerBlock;
cuFuncSetBlockShape(vecAdd, threadsPerBlock, 1, 1);
cuLaunchGrid(VecAdd, blocksPerGrid, 1);
// Copy result from device memory to host memory
// h_C contains the result in host memory
cuMemcpyDtoH(h_C, d_C, size);
// Free device memory
cuMemFree(d_A);
cuMemFree(d_B);
cuMemFree(d_C);
return (0);
}