floating point atomic adds in Fermi
float atomicAdd(float *address, float val)
Intriguing is a mutex also
extern void CUDARTAPI __cudaMutexOperation(int lock);
#define __cudaAtomicOperation(code) \
__cudaMutexOperation(1); \
code \
__cudaMutexOperation(0);
with is used to implement the atomic float add
__device_func__(float __fAtomicAdd(float *address, float val))
{
float old;
__cudaAtomicOperation(
old = *address;
*address = old + val;
)
return old;
}
cuda gl texture interop
==============
get device: cudaGLSetGLDevice (int device)
create tex
register tex: cudaGraphicsGLRegisterImage
setup
while()
{
cudaGraphicsMapResources
get array: cudaGraphicsSubResourceGetMappedArray
use cuda kernel
cudaGraphicsUnmapResources
draw gl
}
new SOs:
Ubuntu 9.04*
Mac OS X 10.6.0*
Mac OS X 10.6.1*
New API Features
--------------------------
- Added the BLAS1 functions:
* cublasZaxpy()
* cublasZcopy()
* cublasZswap()
- Added the BLAS2 functions:
* cublasDtrmv()
* cublasCtrmv()
* cublasCgemv()
* cublasCgeru()
* cublasCgerc()
* cublasZtrmv()
* cublasZgemv()
* cublasZgeru()
* cublasZgerc()
- Added the BLAS3 functions:
* cublasCtrsm()
* cublasCtrmm()
* cublasCsyrk()
* cublasCsymm()
* cublasCherk()
* cublasZtrsm()
* cublasZtrmm()
* cublasZsyrk()
* cublasZsymm()
* cublasZherk()
o Float16 (half) textures are supported in the runtime
- cudaCreateChannelDescHalf family of functions supports it in C++ style
API or proper channel could be crated via cudaCreateChannelDesc in C
style level API
- users should be aware that halves are promoted to floats during
computation and therefore, only floats could be fetched by texture
fetch functions
- users could use intrinsics in device code to convert between fp16 and
fp32 data
o Double3 and double4 vector types are supported in the runtime
- This breaks code when users had already added these themselves.
o One dimensional device-device copies now support streams.
- cudaMemcpyAsync now applies the stream parameter for
cudaMemcpyDeviceToDevice as well
- cuMemcpyDtoDAsync
o Support for ELF binaries
- ELF is generated by default by nvcc. For ptxas or fatbin, the -elf option
is required.
- Cubins are now binary files. Do not assume that they are ASCII text.
o Testing applications for Fermi-readiness
- Setting the env variable CUDA_FORCE_PTX_JIT to 1 will disable all non-PTX
user kernels from being able to load. If your application fails to run,
you are not compiling with PTX. Please see the programming guide for more
information about compiling for different compute capabilities.
o OpenGL texture interoperation
o Batched 2D & 3D transforms are now supported in CUFFT, using the new
cufftPlanMany() API. This is defined in cufft.h, as follows:
cufftResult CUFFTAPI cufftPlanMany(cufftHandle *plan,
int rank,
int *n,
int *inembed, // Unused: pass NULL
int istride, // Unused: pass 1
int idist, // Unused: pass 0
int *onembed, // Unused: pass NULL
int ostride, // Unused: pass 1
int odist, // Unused: pass 0
cufftType type,
int batch);
The arguments are:
*plan - The plan is returned here, as for other cufft calls
rank - The dimensionality of the transform (1, 2 or 3)
*n - An array of size [rank], describing the size of each
dimension
type - Transform type (e.g. CUFFT_C2C), as per other cufft calls
batch - Batch size for this transform
Return values are as for all other cufftPlanXxx functions. Thus to plan
a batch of 1000, 2D, double-precision, complex-to-complex transforms of
size (128, 256), you would do:
cufftHandle *myplan;
cufftPlanMany(myplan, 2, { 128, 256 }, NULL, 1, 0, NULL, 1, 0, CUFFT_Z2Z, 1000);
Note that for CUFFT 3.0, the layout of batched data must be side-by-side
and not interleaved. The inembed, istride, idist, onembed, ostride and
odist parameters are for enabling data windowing and interleaving in a
future version.
New Toolkit Features
--------------------------
o nvcc
- The command line option --host-compilation=C is no more. nvcc emits a
warning and switches back to C++. This option will eventually disappear
altogether
o CUDA GDB known issues:
- Please see the "Known Issues" section in the CUDA_GDB_v3.0.pdf User Manual.
o Windows DLL Naming Conventions
- Each DLL now specifies the machine type, the toolkit version number, and
the build number in its filename.
- For example, cudart32_30_4.dll would be the 32-bit build of 3.0 Cudart
with a build number of 4.
- The build number of the final release will always be greater than the
build number of the beta release.
- The corresponding .lib files do not have any extra naming decoration, so
you can continue linking your applications the same way.
o Separate Library for Runtime Device Emulation
- Cudart has now been split up into two libraries. Cudartemu should be
linked with for device emulation, similar to the way in which
Cublasemu/Cufftemu were previously used.
--------------------------------------------------------------------------------
Bug Fixes
--------------------------------------------------------------------------------
o The asynchronous memcpy routines require the user to pass pinned memory
allocations for any host pointers. In Cuda 2.1, 2.2, and 2.3, no error was
returned if you used non-pinned memory with the NULL stream in some
Host-to-Device memcpy operations. This release adds back the appropriate
error check and returns cudaErrorInvalidValue or CUDA_ERROR_INVALID_VALUE
when an application uses non-pinned memory in such a transfer.
o Both the cudaEventQuery() and cudaStreamQuery() functions have been altered
such that they longer show first-chance exceptions when cudaErrorNotReady
would be returned. This eliminates an issue where users could not turn on
exception debugging in Visual Studio for applications that used these API
calls.
= Known Issues =
CUDA Visual Profiler is non-functional on MacOS in this release. This will be resolved in the production release.
The new cuda-memcheck utility is missing from the CUDA Toolkit packages for Linux 64-bit systems. It is included as a separate package called cuda-memcheck_3.0beta1_linux64.tar.gz and should be installed by:
1. download unpack the archive: tar xfz
2. copy cuda-memcheck to the location you want to be in
3. before running set the LD_LIBRARY_PATH,
If you use bash: export LD_LIBRARY_PATH=/usr/local/cuda/lib:/usr/local/cuda/lib64
OpenCL Visual Profiler for 64-bit Windows and Linux systems is available as a separate package. To install do the following:
For Windows
1. download and unpack openclprof_3.0-beta1_win_64.zip
2. follow steps in OpenCL_Visual_Profiler_Release_Notes.txt under
the top level directory.
For Linux
1. download and unpack openclprof_3.0-beta1_linux_64.tar.gz
2. follow steps in OpenCL_Visual_Profiler_Release_Notes.txt under
the openclprof directory.
cuda-gdb does not support debugging just-in-time (JIT) compiled PTX kernels in this release.
You may notice some image corruption when using OpenGL interop in multi-GPU systems where the GPU used for computation is different than the GPU used for graphics. This will be resolved in the production release.
The RadixSort SDK code sample does not run on Linux32.
The SamplesimpleD3D9Texture does not run on Window 32/64.
Questions i have:
Using DP, cache is hardware?
mac 64 bit?
OpenCL support using two command_queues?:
- Multiple Copy Engine support
- Concurrent Kernel Execution
For fermi usage..
OpenCL and ECC reporting..
no updated docs
ptx 1.5 2.0 docs
Fermi howto?
predication
use host calls
alloc mem in kernel
recursivity
0 comments:
Post a Comment