blob: 14610d8a321ab7f6e1386c6f99b5008120701f30 (
plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
|
/*
* 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 <stdio.h>
#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 2.0f*bar;
}
__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])/2.0;
__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_
|
