Showing posts with label CUDA Basics. Show all posts

How to specify architecture while compiling CUDA program in Visual profiler

How to change command line flag in CUDA

Hello, today I’ll let you know how to change compilation option such as flag in visual profiler. Sometimes it happens we need to specify the architecture of our GPU to run any sample code. For example, if your code contain “atomicAdd()” function then you SM architecture
should be > SM10, although you SM architecture is greater than 1.0 but you can’t Build your project in Visual profiler because of compiler flag is SM10. When you try to do this visual profiler through an error like this;

atomicAdd undefined identifier

In this case you specifically, specify architecture greater than SM10 in compiler flag.

Let us compile our code using command line;

nvcc –arch=sm__30

Same you want to do in your visual profiler. So you need to change this line

nvcc -gencode=arch=compute_10,code=\"sm_30,compute_30\

nvcc -gencode=arch=compute_30,code=\"sm_30,compute_30\

Procedure to achieve this;

Step 1: Open your project properties

Step 2: Go to your GPU option

Step 3: Change GPU architecture value to required SM value, in this case I have change it from SM10 to SM30

Step 4 : You are done

Got Questions?

Feel free to ask me any question by commenting below because I'd be happy to walk you through step by step!

Want to Contact us? Click here

Continuing from Part2

Part3

Further Optimization in histogram CUDA code | Fast implementation of histogram in CUDA

Optimization step 4

HISTOGRAM KERNEL USING SHARED AND GLOBAL MEMORY ATOMICS

Up to this step, our performance on GPU was really bad as compare to CPU version. The one of the possible optimization I though start using shared memory. If reader is not familiar with shared memory concept feel free to follow this link, I explained shared memory with wide details.
So, I’m damn sure all the reader must thought how this optimization comes in my mind?
Well answer is trivial, since when we using global memory, the memory traffic is terrible and even you can use Visual profiler to profile your code to get to know about the performance, or you can build a performance matrix for that, follow this link for that.

Apart from the shared memory there is another way to optimize it, which we have discussed in “Optimization step 3” lunching two dimension blocks and threads. Well, I’m interested to optimize my code first with shared memory than I’ll configure my blocks and grids dimension.

Motivation for Shared memory

Ironically, despite that the atomic operations cause this performance degradation, alleviating the slowdown actually involves using more atomics, not fewer. The core problem was not the use of atomics so much as the fact that thousands of threads were competing for access to a relatively small number of memory addresses.

To address this issue, we will split our histogram computation into two phases.

Phase 1

Computing histogram in shared memory

In phase one, each parallel block will compute a separate histogram of the data that its constituent threads examine. Since each block does this independently, we can compute these histograms in shared memory, saving us the time of sending each write-off chip to DRAM. Doing this does not free us from needing atomic operations, though, since multiple threads within the block can still examine data elements with the same value. However, the fact that only 256 threads will now be competing for 256 addresses will reduce contention from the global version where thousands of threads were competing.

The first phase then involves allocating and zeroing a shared memory buffer to hold each block’s intermediate histogram. Recall from shared memory article, that since the subsequent step will involve reading and modifying this buffer, we need a __syncthreads() call to ensure that every thread’s write has completed before progressing.

__global__ void histo_kernel_optimized5 ( unsigned char *buffer, long size,

unsigned int *histo )

{

__shared__ unsigned int temp[256];

temp[threadIdx.x] = 0;

__syncthreads();

After zeroing the histogram, the next step is remarkably similar to our original GPU histogram. The sole differences here are that we use the shared memory buffer temp[] instead of the global memory buffer histo[] and that we need a subsequent call to __syncthreads() to ensure the last of our writes have been committed.

int i = threadIdx.x + blockIdx.x * blockDim.x;

int offset = blockDim.x * gridDim.x;

while (i < size)

{

atomicAdd( &temp[buffer[i]], 1);

i += offset;

}

__syncthreads();

Phase 2

Merging histograms

The last step in our modified histogram example requires that we merge each block’s temporary histogram into the global buffer histo[].

Suppose we split the input in half and two threads look at different halves and compute separate histograms. If thread A sees byte 0xFC 20 times in the input and thread B sees byte 0xFC 5 times, the byte 0xFC must have appeared 25 times in the input. Likewise, each bin of the final histogram is just the sum of the corresponding bin in thread A’s histogram and thread B’s histogram. This logic extends to any number of threads, so merging every block’s histogram into a single final histogram involves adding each entry in the block’s histogram to the corresponding entry in the final histogram. For all the reasons we’ve seen already, this needs to be done atomically:

atomicAdd( &(histo[threadIdx.x]), temp[threadIdx.x] );

}

Important note

Since we have decided to use 256 threads and have 256 histogram bins, each thread atomically adds a single bin to the final histogram’s total. If these numbers didn’t match, this phase would be more complicated. Note that we have no guarantees about what order the blocks add their values to the final histogram, but since integer addition is commutative, we will always get the same answer provided that the additions occur atomically.

And with this, our two phase histogram computation kernel is complete. Here it is from start to finish:

__global__ void histo_kernel_optimized5( unsigned char *buffer, long size,

unsigned int *histo )

{

__shared__ unsigned int temp[256];

temp[threadIdx.x] = 0;

__syncthreads();

int i = threadIdx.x + blockIdx.x * blockDim.x;

int offset = blockDim.x * gridDim.x;

while (i < size)

{

atomicAdd( &temp[buffer[i]], 1);

i += offset;

}

__syncthreads();

atomicAdd( &(histo[threadIdx.x]), temp[threadIdx.x] );

}

This version of our histogram example improves dramatically over the previous GPU version. Adding the shared memory component drops our running time on a GeForce GTX 285 to 0.057 seconds. Not only is this significantly better than the version that used global memory atomics only, but this beats our original CPU implementation by an order of magnitude (from 0.416 seconds to 0.057 seconds). This improvement represents greater than a sevenfold boost in speed over the CPU version. So despite the early setback in adapting the histogram to a GPU implementation, our version that uses both shared and global atomics should be considered a success.

Can we do more faster

Yes we can do, after profiling above code, I found it has bank conflict. So if we are able to remove bank conflict then time can reduce by factor 2 or 4 or 8 or even more.
I left this task now for reader,

I’ll present bank conflict free histogram computation in subsequent article.

Now what about competition,

In that online competition, number of bins was 1024 and Size = 1000*1024 = 1024000 elements.

my naïve implementation (optimization 3) was performed well and time was 16.xx ms, with optimization step 4 implementation, time was 8.xx ms ( 2x faster ) and with the optimization step 5 implementation, time was 1.xx ms, it means 8 time faster than global memory version and 16 time faster than naïve implementation.

My kernel invocation was look like this;

// with naïve implementation

naive_kernel <<<blocks*8,1024>>>( d_vals, d_histo, numElems );

//with optimized implementaion

OptimizedStep5 <<<blocks*8,1024, 1024*sizeof(int)>>>( d_vals, d_histo,

numElems );

You may noticed, I lunched 1024 threads with dynamic shared memory allocation (look third argument of kernel invocation) since number of bins was 1024.

Further Problem statement

If these numbers of bins and threads in a block does not match than this problem will be more complicate, like as above, number of bins was 1024 and let say I can lunch only 512 threads in a block then how to solve this problem?

I've cover this approach in Part4 of this article.

All In one CUDA

CUDA syntax

Source code is in .cu files, which contain mixture of host (CPU) and device (GPU) code.

Declaring functions

`__global__`	declares kernel, which is called on host and executed on device
`__device__`	declares device function, which is called and executed on device
`__host__`	declares host function, which is called and executed on host
`__noinline__`	to avoid inlining
`__forceinline__`	to force inlining

Declaring variables

`__device__`	declares device variable in global memory, accessible from all threads, with lifetime of application
`__constant__`	declares device variable in constant memory, accessible from all threads, with lifetime of application
`__shared__`	declares device varibale in block's shared memory, accessible from all threads within a block, with lifetime of block
`__restrict__`	standard C definition that pointers are not aliased

Types

Most routines return an error code of type cudaError_t.

Vector types

char1, uchar1, short1, ushort1, int1, uint1, long1, ulong1, float1
char2, uchar2, short2, ushort2, int2, uint2, long2, ulong2, float2
char3, uchar3, short3, ushort3, int3, uint3, long3, ulong3, float3
char4, uchar4, short4, ushort4, int4, uint4, long4, ulong4, float4

longlong1, ulonglong1, double1
longlong2, ulonglong2, double2

dim3

Components are accessible as variable.x, variable.y, variable.z, variable.w.
Constructor is make_<type>( x, ... ), for example:

float2 xx = make_float2( 1., 2. );

dim3 can take 1, 2, or 3 argumetns:

dim3 blocks1D( 5       );
dim3 blocks2D( 5, 5    );
dim3 blocks3D( 5, 5, 5 );

Pre-defined variables

`dim3 gridDim`	dimensions of grid
`dim3 blockDim`	dimensions of block
`uint3 blockIdx`	block index within grid
`uint3 threadIdx`	thread index within block
`int warpSize`	number of threads in warp

Kernel invocation

__global__ void kernel( ... ) { ... }

dim3 blocks( nx, ny, nz );           // cuda 1.x has 1D and 2D grids, cuda 2.x adds 3D grids
dim3 threadsPerBlock( mx, my, mz );  // cuda 1.x has 1D, 2D, and 3D blocks

kernel<<< blocks, threadsPerBlock >>>( ... );

Thread management

`__threadfence_block();`	wait until memory accesses are visible to block
`__threadfence();`	wait until memory accesses are visible to block and device
`__threadfence_system();`	wait until memory accesses are visible to block and device and host (2.x)
`__syncthreads();`	wait until all threads reach sync

Memory management

__device__ float* pointer;
cudaMalloc( &pointer, size );
cudaFree( pointer );

// direction is one of cudaMemcpyHostToDevice or cudaMemcpyDeviceToHost
cudaMemcpy( dst_pointer, src_pointer, size, direction );

__constant__ float dev_data[n];
float host_data[n];
cudaMemcpyToSymbol  ( dev_data,  host_data, sizeof(host_data) );  // dev_data  = host_data
cudaMemcpyFromSymbol( host_data, dev_data,  sizeof(host_data) );  // host_data = dev_data

Also, malloc and free work inside a kernel (2.x), but memory allocated in a kernel must be deallocated in a kernel (not the host). It can be freed in a different kernel, though.

Atomic functions

old = atomicAdd ( &addr, value );  // old = *addr;  *addr += value
old = atomicSub ( &addr, value );  // old = *addr;  *addr –= value
old = atomicExch( &addr, value );  // old = *addr;  *addr  = value

old = atomicMin ( &addr, value );  // old = *addr;  *addr = min( old, value )
old = atomicMax ( &addr, value );  // old = *addr;  *addr = max( old, value )

// increment up to value, then reset to 0  
// decrement down to 0, then reset to value
old = atomicInc ( &addr, value );  // old = *addr;  *addr = ((old >= value) ? 0 : old+1 )
old = atomicDec ( &addr, value );  // old = *addr;  *addr = ((old == 0) or (old > val) ? val : old–1 )

old = atomicAnd ( &addr, value );  // old = *addr;  *addr &= value
old = atomicOr  ( &addr, value );  // old = *addr;  *addr |= value
old = atomicXor ( &addr, value );  // old = *addr;  *addr ^= value

// compare-and-store
old = atomicCAS ( &addr, compare, value );  // old = *addr;  *addr = ((old == compare) ? value : old)

Warp vote

int __all   ( predicate );
int __any   ( predicate );
int __ballot( predicate );  // nth thread sets nth bit to predicate

Timer

wall clock cycle counter

clock_t clock();

Texture

can also return float2 or float4, depending on texRef.

// integer index
float tex1Dfetch( texRef, ix );

// float index
float tex1D( texRef, x       );
float tex2D( texRef, x, y    );
float tex3D( texRef, x, y, z );

float tex1DLayered( texRef, x    );
float tex2DLayered( texRef, x, y );

Low-level Driver API

#include <cuda.h>

CUdevice dev;
CUdevprop properties;
char name[n];
int major, minor;
size_t bytes;

cuInit( 0 );  // takes flags for future use
cuDeviceGetCount         ( &cnt );
cuDeviceGet              ( &dev, index );
cuDeviceGetName          ( name, sizeof(name), dev );
cuDeviceComputeCapability( &major, &minor,     dev );
cuDeviceTotalMem         ( &bytes,             dev );
cuDeviceGetProperties    ( &properties,        dev );  // max threads, etc.

cuBLAS

Matrices are column-major. Indices are 1-based; this affects result of i<t>amax and i<t>amin.

#include <cublas_v2.h>

cublasHandle_t handle;
cudaStream_t   stream;

cublasCreate( &handle );
cublasDestroy( handle );
cublasGetVersion( handle, &version );
cublasSetStream( handle,  stream );
cublasGetStream( handle, &stream );
cublasSetPointerMode( handle,  mode );
cublasGetPointerMode( handle, &mode );

// copy x => y
cublasSetVector     ( n, elemSize, x_src_host, incx, y_dst_dev,  incy );
cublasGetVector     ( n, elemSize, x_src_dev,  incx, y_dst_host, incy );
cublasSetVectorAsync( n, elemSize, x_src_host, incx, y_dst_dev,  incy, stream );
cublasGetVectorAsync( n, elemSize, x_src_dev,  incx, y_dst_host, incy, stream );

// copy A => B
cublasSetMatrix     ( rows, cols, elemSize, A_src_host, lda, B_dst_dev,  ldb );
cublasGetMatrix     ( rows, cols, elemSize, A_src_dev,  lda, B_dst_host, ldb );
cublasSetMatrixAsync( rows, cols, elemSize, A_src_host, lda, B_dst_dev,  ldb, stream );
cublasGetMatrixAsync( rows, cols, elemSize, A_src_dev,  lda, B_dst_host, ldb, stream );

Constants

argument	constants	description (Fortran letter)
`trans`	`CUBLAS_OP_N`	non-transposed ('N')
	`CUBLAS_OP_T`	transposed ('T')
	`CUBLAS_OP_C`	conjugate transposed ('C')

`uplo`	`CUBLAS_FILL_MODE_LOWER`	lower part filled ('L')
	`CUBLAS_FILL_MODE_UPPER`	upper part filled ('U')

`side`	`CUBLAS_SIDE_LEFT`	matrix on left ('L')
	`CUBLAS_SIDE_RIGHT`	matrix on right ('R')

`mode`	`CUBLAS_POINTER_MODE_HOST`	alpha and beta scalars passed on host
	`CUBLAS_POINTER_MODE_DEVICE`	alpha and beta scalars passed on device

BLAS functions have cublas prefix and first letter of usual BLAS function name is capitalized. Arguments are the same as standard BLAS, with these exceptions:

All functions add handle as first argument.
All functions return cublasStatus_t error code.
Constants alpha and beta are passed by pointer. All other scalars (n, incx, etc.) are bassed by value.
Functions that return a value, such as ddot, add result as last argument, and save value to result.
Constants are given in table above, instead of using characters.

Examples:

cublasDdot ( handle, n, x, incx, y, incy, &result );  // result = ddot( n, x, incx, y, incy );
cublasDaxpy( handle, n, &alpha, x, incx, y, incy );   // daxpy( n, alpha, x, incx, y, incy );

Compiler

nvcc, often found in /usr/local/cuda/bin

Defines __CUDACC__

Flags common with cc

Short flag	Long flag	Output or Description
`-c`	`--compile`	.o object file
`-E`	`--preprocess`	on standard output
`-M`	`--generate-dependencies`	on standard output
`-o file`	`--output-file file`
`-I directory`	`--include-path directory`	header search path
`-L directory`	`--library-path directory`	library search path
`-l lib`	`--library lib`	link with library
`-lib`		generate library
`-shared`		generate shared library
`-pg`	`--profile`	for gprof
`-g level`	`--debug level`
`-G`	`--device-debug`
`-O level`	`--optimize level`

Undocumented (but in sample makefiles)
`-m64`		compile x86_64 host CPU code

Flags specific to nvcc

`-v`	list compilation commands as they are executed
`-dryrun`	list compilation commands, without executing
`-keep`	saves intermediate files (e.g., pre-processed) for debugging
`-clean`	removes output files (with same exact compiler options)
`-arch=<compute_xy>`	generate PTX for capability x.y
`-code=<sm_xy>`	generate binary for capability x.y, by default same as `-arch`
`-gencode arch=...,code=...`	same as `-arch` and `-code`, but may be repeated

Argumenents for `-arch` and `-code`

It makes most sense (to me) to give -arch a virtual architecture and -code a real architecture, though both flags accept both virtual and real architectures (at times).

	Virtual architecture	Real architecture	Features
Tesla	`compute_10`	`sm_10`	Basic features
	`compute_11`	`sm_11`	+ atomic memory ops on global memory
	`compute_12`	`sm_12`	+ atomic memory ops on shared memory + vote instructions
	`compute_13`	`sm_13`	+ double precision
Fermi	`compute_20`	`sm_20`	+ Fermi

Some hardware constraints

	1.x	2.x
max x- or y-dimension of block	512	1024
max z-dimension of block	64	64
max threads per block	512	1024
warp size	32	32
max blocks per MP	8	8
max warps per MP	32	48
max threads per MP	1024	1536
max 32-bit registers per MP	16k	32k
max shared memory per MP	16 KB	48 KB
shared memory banks	16	32
local memory per thread	16 KB	512 KB
const memory	64 KB	64 KB
const cache	8 KB	8 KB
texture cache	8 KB	8 KB

Got Questions? Feel free to ask me any question because I'd be happy to walk you through step by step! References For Contact us….. Click on Contact us Tab

So many visitors have asked me that,

“How to install NVidia Nsight Visual Studio plugin in Visual Studio for Debugging CUDA Code?”

“How to configure NVidia Nsight in Visual Studio for debugging CUDA code?”

Today I’m going to teach you a step by step guide to resolve above query. Please follow below steps for correct installation and configuration afterword’s, I’ll show you a video that will let you know How to use Nvidia Nsight Visual studio plugin for debugging CUDA code.

Step1: Get to know about requirement, follow below link and check that your Card and visual studio version met the requirement or not. Nsight Visual Studio Edition Requirements.

Step 2: Get to know Nsight Visual Studio Edition Supported GPUs, follow below link Nsight Visual Studio Edition Supported GPUs

Step3: To download Nsight Visual studio plugin you need to register at Developer website of NVidia. If you have an account already then download from here otherwise register at here (After registration, they send you a download link), download your Nsight Visual studio plugin setup and save in your Hard disk.

Step4: Start installation by double clicking on Setup of Nsight, Installation will start automatically. Be careful at this step, this is most important step at which you need to concentrate more. When you start installing check does any requirement doesn’t met, if yes then follow this link Install Nsight Visual Studio Edition

Completed all steps? Yes, Hurry you install Nsight perfectly (please verify it). Now it’s time to learn how to use it, here is the video which describe you how to use it.

Debugging CUDA C using NVIDIA Parallel Nsight and Visual Studio 2008

Debugging Graphics using NVIDIA Parallel Nsight and Visual Studio

Got Questions?

Feel free to ask me any question because I'd be happy to walk you through step by step!

References and External links

Nsight Visual Studio Edition Requirements.

Nsight Visual Studio Edition Supported GPUs

Install Nsight Visual Studio Edition

For Contact us….. Click on Contact us Tab

Handling CUDA error messages

From Visualization

The following piece of code, in the cudasafe routine, can be used to handle CUDA error messages.

1. #include <iostream>

2. #include <cuda.h>

4. void cudasafe( cudaError_t error, char* message)

5. {

6. if(error!=cudaSuccess) { fprintf(stderr,"ERROR: %s : %i\n",message,error); exit(-1); }

7. }

9. int main() {

10. float *a_d; // pointers to device memory; a.k.a. GPU

11. int block_size, block_no, n=10;

12.

13.// allocate arrays on device

14. cudasafe( cudaMalloc((void **)&a_d,n*n*sizeof(float)), "cudaMalloc" );

15. block_size=22;

16. dim3 dimBlock( block_size, block_size );

17. dim3 dimGrid( ceil(float(N)/float(dimBlock.x)), ceil(float(N)/float(dimBlock.y)) );

18.

19. cudasafe( cudaFree(a_d), "cudaFree" );

20.

21. return(0);

22.}

See the Example 1 for another way to check for error messages with CUDA.

Common Errors

error: a host function call can not be configured simply means that you tried to call a routine as if it was a kernel to be executed on the device, but you forgot to put __global__ in front of that routine.

Invalid Configuration Argument - This error means that the dimension of either the specified grid of blocks (dimGrid) , or number of threads in a block (dimBlock), is incorrect. In such a case, the dimension is either zero or the dimension is larger than it should be. This error will only occur if you dynamically determine the dimensions.

Too Many Resources Requested for Launch - This error means that the number of registers available on the multiprocessor is being exceeded. Reduce the number of threads per block to solve the problem.

Unspecified launch failure - This error means that CUDA does not know what the problem was. This is the worst error to get because you do not know where to look to correct the error. One way to look at this error message is to mentally translate it to "segmentation fault" for the host code.

Example 1

Suppose you used the following piece of code in your program to check for error messages.

1. void Check_CUDA_Error(const char *message)

2. {

3. cudaError_t error = cudaGetLastError();

4. if(error!=cudaSuccess) {

5. fprintf(stderr,"ERROR: %s: %s\n", message, cudaGetErrorString(error) );

6. exit(-1);

7. }

8. }

9. int main(int argc, char** argv))

10.{

11. :

12. :

13. block_size=23;

14. dim3 dimBlock(block_size,block_size);

15. dim3 dimGrid( ceil(float(N)/float(dimBlock.x)), ceil(float(N)/float(dimBlock.y)));

16. assign_d<<<dimGrid, dimBlock>>>(a_d,N);

17. Check_CUDA_Error("Kernel Execution Failed!");

18. :

19. :

20. return 0;

21.}

This piece of code would fail without a warning as to the cause. Remember, once you launch the kernel, it operates asynchronously with the CPU. The kernel would fail, and not tell you, but the CPU would continue to compute whatever was left in the program. By checking the error message, you could see that the kernel failed with Invalid Configuration Argument. In this case, we know number of threads in the block is not zero. However, there are 529 threads in the block, which exceeds the capability of the GPU, which was shown in the Getting information about the GPU tutorial to be 512. By reducing the number of threads down to 22 per side of the block (or 484 threads total in the block), the code will run correctly.

Example 2

Suppose you took the code from the Laplace Solver Program tutorial, and modified it so that instead of:

1. if(i>0 && i<N-1 && j>0 && j<N-1) { B[index] = 0.25*( A[index1] + A[index2] + A[index3] + A[index4] ); }

you now had a new array, a masking array. This masking array is set to zero on the boundaries of the array, and one on the interior. This way the interior is computed, and the boundary conditions are left alone.

1. if(mask[index]) { B[index] = 0.25*( A[index1] + A[index2] + A[index3] + A[index4] ); }

However, when you run the code, you occasionally get the dreaded unspecified launch failure error. Sometimes when you run the code it works fine; sometimes it fails. The problem is that you are accessing an array out of bounds, which is giving you the error. When the program is executed, a number of threads are created. These threads are grouped together in thread blocks. Suppose you say you want 16 threads per block, and the grid on which you are solving the Laplace Equation is 45 x 45. The grid has 2,025 points. Each block has 256 threads. Dividing the size of the grid by the number of threads per block means that you will need 7.9 blocks. Of course, you cannot have a partial block, so the number is rounded up to 8 blocks. That means that you have 2,048 threads, while you need only 2,025. (Really you only need 1,936 threads since you have boundary conditions where no computation takes place.) The extra threads are unused in the first code block. However, in the second code block, with the masking array, those extra threads will be accessing the mask array beyond the bounds of the array. The result is non-deterministic. Sometimes it may succeed ; sometimes it may fail with the error unspecified launch failure.

Got Questions?

Feel free to ask me any question because I'd be happy to walk you through step by step!

For Contact us….. Click on Contact us Tab

CUDA Programming

Prefer Your Language

Search This Blog

How to specify architecture while compiling CUDA program in Visual profiler How to change command line flag in CUDA

Further Optimization in histogram CUDA code | Fast implementation of histogram in CUDA

CUDA complete | Complete reference on CUDA

CUDA syntax

Declaring functions

Declaring variables

Types

Vector types

Pre-defined variables

Kernel invocation

Thread management

Memory management

Atomic functions

Warp vote

Timer

Texture

Low-level Driver API

cuBLAS

Constants

Compiler

Flags common with cc

Flags specific to nvcc

Argumenents for -arch and -code

Some hardware constraints

Installing NVidia Nsight Visual studio plugin for Debugging CUDA Kernel Code

Step 2: Get to know Nsight Visual Studio Edition Supported GPUs, follow below link Nsight Visual Studio Edition Supported GPUs

Completed all steps? Yes, Hurry you install Nsight perfectly (please verify it). Now it’s time to learn how to use it, here is the video which describe you how to use it.

Debugging CUDA C using NVIDIA Parallel Nsight and Visual Studio 2008

Debugging Graphics using NVIDIA Parallel Nsight and Visual Studio

Handling CUDA error messages

Recent Post

About Me

Subscribe To

Total Pageviews

Labels

Blog Archive

Labels

Like Us

Cloud

Admin

Argumenents for `-arch` and `-code`