From 11b308e7363e937966b035b4891db32b4eece3bf Mon Sep 17 00:00:00 2001 From: Tor Aamodt Date: Fri, 1 Oct 2010 08:55:28 -0800 Subject: integrating recent changes from fermi-test into fermi (i'll use "fermi" for more disruptive changes to the pipeline model such as updating the MSHRs and getting rid of the warp tracker, ripping out DWF, etc...) [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7805] --- benchmarks/CUDA/template/Makefile | 54 ----------- benchmarks/CUDA/template/template.cu | 136 ---------------------------- benchmarks/CUDA/template/template_gold.cpp | 37 -------- benchmarks/CUDA/template/template_kernel.cu | 66 -------------- 4 files changed, 293 deletions(-) delete mode 100644 benchmarks/CUDA/template/Makefile delete mode 100644 benchmarks/CUDA/template/template.cu delete mode 100644 benchmarks/CUDA/template/template_gold.cpp delete mode 100644 benchmarks/CUDA/template/template_kernel.cu (limited to 'benchmarks/CUDA/template') diff --git a/benchmarks/CUDA/template/Makefile b/benchmarks/CUDA/template/Makefile deleted file mode 100644 index a36bce5..0000000 --- a/benchmarks/CUDA/template/Makefile +++ /dev/null @@ -1,54 +0,0 @@ -################################################################################ -# -# Copyright 1993-2006 NVIDIA Corporation. All rights reserved. -# -# NOTICE TO USER: -# -# This source code is subject to NVIDIA ownership rights under U.S. and -# international Copyright laws. -# -# NVIDIA MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THIS SOURCE -# CODE FOR ANY PURPOSE. IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR -# IMPLIED WARRANTY OF ANY KIND. NVIDIA DISCLAIMS ALL WARRANTIES WITH -# REGARD TO THIS SOURCE CODE, INCLUDING ALL IMPLIED WARRANTIES OF -# MERCHANTABILITY, NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE. -# IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL, -# OR CONSEQUENTIAL DAMAGES, OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS -# OF USE, DATA OR PROFITS, WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE -# OR OTHER TORTIOUS ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE -# OR PERFORMANCE OF THIS SOURCE CODE. -# -# U.S. Government End Users. This source code is a "commercial item" as -# that term is defined at 48 C.F.R. 2.101 (OCT 1995), consisting of -# "commercial computer software" and "commercial computer software -# documentation" as such terms are used in 48 C.F.R. 12.212 (SEPT 1995) -# and is provided to the U.S. Government only as a commercial end item. -# Consistent with 48 C.F.R.12.212 and 48 C.F.R. 227.7202-1 through -# 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the -# source code with only those rights set forth herein. -# -################################################################################ -# -# Build script for project -# -################################################################################ - -# Add source files here -EXECUTABLE := template -# CUDA source files (compiled with cudacc) -CUFILES := template.cu -# CUDA dependency files -CU_DEPS := \ - template_kernel.cu \ - -# C/C++ source files (compiled with gcc / c++) -CCFILES := \ - template_gold.cpp \ - - -################################################################################ -# Rules and targets - -ROOTDIR=$(NVIDIA_CUDA_SDK_LOCATION)/src -BINDIR=. -include $(NVIDIA_CUDA_SDK_LOCATION)/common/common.mk diff --git a/benchmarks/CUDA/template/template.cu b/benchmarks/CUDA/template/template.cu deleted file mode 100644 index 8a44200..0000000 --- a/benchmarks/CUDA/template/template.cu +++ /dev/null @@ -1,136 +0,0 @@ -/* - * Copyright 1993-2010 NVIDIA Corporation. All rights reserved. - * - * NVIDIA Corporation and its licensors retain all intellectual property and - * proprietary rights in and to this software and related documentation. - * Any use, reproduction, disclosure, or distribution of this software - * and related documentation without an express license agreement from - * NVIDIA Corporation is strictly prohibited. - * - * Please refer to the applicable NVIDIA end user license agreement (EULA) - * associated with this source code for terms and conditions that govern - * your use of this NVIDIA software. - * - */ - -/* Template project which demonstrates the basics on how to setup a project -* example application. -* Host code. -*/ - -// includes, system -#include -#include -#include -#include - -// includes, project -#include - -// includes, kernels -#include - -//////////////////////////////////////////////////////////////////////////////// -// declaration, forward -void runTest( int argc, char** argv); - -extern "C" -void computeGold( float* reference, float* idata, const unsigned int len); - -//////////////////////////////////////////////////////////////////////////////// -// Program main -//////////////////////////////////////////////////////////////////////////////// -int -main( int argc, char** argv) -{ - runTest( argc, argv); - - cutilExit(argc, argv); -} - -//////////////////////////////////////////////////////////////////////////////// -//! Run a simple test for CUDA -//////////////////////////////////////////////////////////////////////////////// -void -runTest( int argc, char** argv) -{ - // use command-line specified CUDA device, otherwise use device with highest Gflops/s - if( cutCheckCmdLineFlag(argc, (const char**)argv, "device") ) - cutilDeviceInit(argc, argv); - else - cudaSetDevice( cutGetMaxGflopsDeviceId() ); - - unsigned int timer = 0; - cutilCheckError( cutCreateTimer( &timer)); - cutilCheckError( cutStartTimer( timer)); - - unsigned int num_threads = 32; - unsigned int mem_size = sizeof( float) * num_threads; - - // allocate host memory - float* h_idata = (float*) malloc( mem_size); - // initalize the memory - for( unsigned int i = 0; i < num_threads; ++i) - { - h_idata[i] = (float) i; - } - - // allocate device memory - float* d_idata; - cutilSafeCall( cudaMalloc( (void**) &d_idata, mem_size)); - // copy host memory to device - cutilSafeCall( cudaMemcpy( d_idata, h_idata, mem_size, - cudaMemcpyHostToDevice) ); - - // allocate device memory for result - float* d_odata; - cutilSafeCall( cudaMalloc( (void**) &d_odata, mem_size)); - - // setup execution parameters - dim3 grid( 1, 1, 1); - dim3 threads( num_threads, 1, 1); - - // execute the kernel - testKernel<<< grid, threads, mem_size >>>( d_idata, d_odata); - - // check if kernel execution generated and error - cutilCheckMsg("Kernel execution failed"); - - // allocate mem for the result on host side - float* h_odata = (float*) malloc( mem_size); - // copy result from device to host - cutilSafeCall( cudaMemcpy( h_odata, d_odata, sizeof( float) * num_threads, - cudaMemcpyDeviceToHost) ); - - cutilCheckError( cutStopTimer( timer)); - printf( "Processing time: %f (ms)\n", cutGetTimerValue( timer)); - cutilCheckError( cutDeleteTimer( timer)); - - // compute reference solution - float* reference = (float*) malloc( mem_size); - computeGold( reference, h_idata, num_threads); - - // check result - if( cutCheckCmdLineFlag( argc, (const char**) argv, "regression")) - { - // write file for regression test - cutilCheckError( cutWriteFilef( "./data/regression.dat", - h_odata, num_threads, 0.0)); - } - else - { - // custom output handling when no regression test running - // in this case check if the result is equivalent to the expected soluion - CUTBoolean res = cutComparef( reference, h_odata, num_threads); - printf( "%s\n", (1 == res) ? "PASSED" : "FAILED"); - } - - // cleanup memory - free( h_idata); - free( h_odata); - free( reference); - cutilSafeCall(cudaFree(d_idata)); - cutilSafeCall(cudaFree(d_odata)); - - cudaThreadExit(); -} diff --git a/benchmarks/CUDA/template/template_gold.cpp b/benchmarks/CUDA/template/template_gold.cpp deleted file mode 100644 index 92a8045..0000000 --- a/benchmarks/CUDA/template/template_gold.cpp +++ /dev/null @@ -1,37 +0,0 @@ -/* - * Copyright 1993-2010 NVIDIA Corporation. All rights reserved. - * - * NVIDIA Corporation and its licensors retain all intellectual property and - * proprietary rights in and to this software and related documentation. - * Any use, reproduction, disclosure, or distribution of this software - * and related documentation without an express license agreement from - * NVIDIA Corporation is strictly prohibited. - * - * Please refer to the applicable NVIDIA end user license agreement (EULA) - * associated with this source code for terms and conditions that govern - * your use of this NVIDIA software. - * - */ - -//////////////////////////////////////////////////////////////////////////////// -// export C interface -extern "C" -void computeGold( float* reference, float* idata, const unsigned int len); - -//////////////////////////////////////////////////////////////////////////////// -//! Compute reference data set -//! Each element is multiplied with the number of threads / array length -//! @param reference reference data, computed but preallocated -//! @param idata input data as provided to device -//! @param len number of elements in reference / idata -//////////////////////////////////////////////////////////////////////////////// -void -computeGold( float* reference, float* idata, const unsigned int len) -{ - const float f_len = static_cast( len); - for( unsigned int i = 0; i < len; ++i) - { - reference[i] = idata[i] * f_len; - } -} - diff --git a/benchmarks/CUDA/template/template_kernel.cu b/benchmarks/CUDA/template/template_kernel.cu deleted file mode 100644 index adac694..0000000 --- a/benchmarks/CUDA/template/template_kernel.cu +++ /dev/null @@ -1,66 +0,0 @@ -/* - * Copyright 1993-2010 NVIDIA Corporation. All rights reserved. - * - * NVIDIA Corporation and its licensors retain all intellectual property and - * proprietary rights in and to this software and related documentation. - * Any use, reproduction, disclosure, or distribution of this software - * and related documentation without an express license agreement from - * NVIDIA Corporation is strictly prohibited. - * - * Please refer to the applicable NVIDIA end user license agreement (EULA) - * associated with this source code for terms and conditions that govern - * your use of this NVIDIA software. - * - */ - -/* Template project which demonstrates the basics on how to setup a project - * example application. - * Device code. - */ - -#ifndef _TEMPLATE_KERNEL_H_ -#define _TEMPLATE_KERNEL_H_ - -#include - -#define SDATA( index) cutilBankChecker(sdata, index) - -//////////////////////////////////////////////////////////////////////////////// -//! Simple test kernel for device functionality -//! @param g_idata input data in global memory -//! @param g_odata output data in global memory -//////////////////////////////////////////////////////////////////////////////// - -__device__ float foo( float bar ) -{ - return bar+1.0f; -} - -__global__ void -testKernel( float* g_idata, float* g_odata) -{ - // shared memory - // the size is determined by the host application - extern __shared__ float sdata[]; - - // access thread id - const unsigned int tid = threadIdx.x; - // access number of threads in this block - const unsigned int num_threads = blockDim.x; - - // read in input data from global memory - // use the bank checker macro to check for bank conflicts during host - // emulation - SDATA(tid) = foo(g_idata[tid])-1.0f; - __syncthreads(); - printf("thread tid=%u reads %f from g_idata[]\n", tid, sdata[tid]); - - // perform some computations - SDATA(tid) = (float) num_threads * SDATA( tid); - __syncthreads(); - - // write data to global memory - g_odata[tid] = SDATA(tid); -} - -#endif // #ifndef _TEMPLATE_KERNEL_H_ -- cgit v1.3