CUFFT_INVALID_PLAN after loading other ptx code with cuModuleLoadDataEx

Sadjad · June 29, 2016, 9:23am

Hello,

I use cuFFT in my application but also some other code that I have compiled into ptx code. After creating the forward transform plan for the fft, I load the ptx code using cuModuleLoadDataEx. If I now call cufftExecR2C with the handle to the forward plan I’ve created before, the function returns CUFFT_INVALID_PLAN.
If I do not load the ptx code, the function succeeds.
I’m a bit confused and I am all new to GPU programming. Is it not possible to load other kernels after creating a fft plan?

Cheers!

Robert_Crovella · June 29, 2016, 1:37pm

It’s possible. You may be doing something wrong. It’s impossible to say without a complete test case.

This may be of interest:

https://devtalk.nvidia.com/default/topic/863142/problems-when-mix-using-cuda-runtime-api-and-cuda-driver-api/?offset=1

Here’s a fully worked example, a modification of the vectorAddDrv CUDA sample code, that adds a cufft plan operation before a cuModuleLoadDataEx call, and a cufft exec operation after that cuModuleLoadDataEx call, successfully.

vectorAddDrv.cpp:

/*
 * Copyright 1993-2015 NVIDIA Corporation.  All rights reserved.
 *
 * Please refer to the NVIDIA end user license agreement (EULA) associated
 * with this source code for terms and conditions that govern your use of
 * this software. Any use, reproduction, disclosure, or distribution of
 * this software and related documentation outside the terms of the EULA
 * is strictly prohibited.
 *
 */

/* Vector addition: C = A + B.
 *
 * This sample is a very basic sample that implements element by element
 * vector addition. It is the same as the sample illustrating Chapter 3
 * of the programming guide with some additions like error checking.
 *
 */

// Includes
#include <stdio.h>
#include <string.h>
#include <iostream>
#include <cstring>

// includes, project
#include <helper_functions.h>
#include <helper_cuda.h>

// includes, CUDA
#include <cuda.h>
#include <cudaProfiler.h>
#include <builtin_types.h>
#include <drvapi_error_string.h>
#include <cuda_runtime_api.h>
#include <cufft.h>

using namespace std;

// Variables
CUdevice cuDevice;
CUcontext cuContext;
CUmodule cuModule;
CUfunction vecAdd_kernel;
float *h_A;
float *h_B;
float *h_C;
CUdeviceptr d_A;
CUdeviceptr d_B;
CUdeviceptr d_C;
bool noprompt = false;

// Functions
void Cleanup(bool);
CUresult CleanupNoFailure();
void RandomInit(float *, int);
bool findModulePath(const char *, string &, char **, string &);
void ParseArguments(int, char **);

int *pArgc = NULL;
char **pArgv = NULL;

//define input ptx file for different platforms
#if defined(_WIN64) || defined(__LP64__)
#define PTX_FILE "vectorAdd_kernel64.ptx"
#else
#define PTX_FILE "vectorAdd_kernel32.ptx"
#endif

////////////////////////////////////////////////////////////////////////////////
// These are CUDA Helper functions

// This will output the proper CUDA error strings in the event that a CUDA host call returns an error
#define checkCudaErrors(err)  __checkCudaErrors (err, __FILE__, __LINE__)

// These are the inline versions for all of the SDK helper functions
inline void __checkCudaErrors(CUresult err, const char *file, const int line)
{
    if (CUDA_SUCCESS != err)
    {
        fprintf(stderr, "checkCudaErrors() Driver API error = %04d \"%s\" from file <%s>, line %i.\n",
                err, getCudaDrvErrorString(err), file, line);
        exit(EXIT_FAILURE);
    }
}

inline int cudaDeviceInit(int ARGC, char **ARGV)
{
    int cuDevice = 0;
    int deviceCount = 0;
    CUresult err = cuInit(0);

    if (CUDA_SUCCESS == err)
    {
        checkCudaErrors(cuDeviceGetCount(&deviceCount));
    }

    if (deviceCount == 0)
    {
        fprintf(stderr, "cudaDeviceInit error: no devices supporting CUDA\n");
        exit(EXIT_FAILURE);
    }

    int dev = 0;
    dev = getCmdLineArgumentInt(ARGC, (const char **) ARGV, "device=");

    if (dev < 0)
    {
        dev = 0;
    }

    if (dev > deviceCount-1)
    {
        fprintf(stderr, "\n");
        fprintf(stderr, ">> %d CUDA capable GPU device(s) detected. <<\n", deviceCount);
        fprintf(stderr, ">> cudaDeviceInit (-device=%d) is not a valid GPU device. <<\n", dev);
        fprintf(stderr, "\n");
        return -dev;
    }

    checkCudaErrors(cuDeviceGet(&cuDevice, dev));
    char name[100];
    cuDeviceGetName(name, 100, cuDevice);

    if (checkCmdLineFlag(ARGC, (const char **) ARGV, "quiet") == false)
    {
        printf("> Using CUDA Device [%d]: %s\n", dev, name);
    }

    return dev;
}

// This function returns the best GPU based on performance
inline int getMaxGflopsDeviceId()
{
    CUdevice current_device = 0, max_perf_device = 0;
    int device_count     = 0, sm_per_multiproc = 0;
    int max_compute_perf = 0, best_SM_arch     = 0;
    int major = 0, minor = 0, multiProcessorCount, clockRate;

    cuInit(0);
    checkCudaErrors(cuDeviceGetCount(&device_count));

    // Find the best major SM Architecture GPU device
    while (current_device < device_count)
    {
        checkCudaErrors(cuDeviceComputeCapability(&major, &minor, current_device));

        if (major > 0 && major < 9999)
        {
            best_SM_arch = MAX(best_SM_arch, major);
        }

        current_device++;
    }

    // Find the best CUDA capable GPU device
    current_device = 0;

    while (current_device < device_count)
    {
        checkCudaErrors(cuDeviceGetAttribute(&multiProcessorCount,
                                             CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT,
                                             current_device));
        checkCudaErrors(cuDeviceGetAttribute(&clockRate,
                                             CU_DEVICE_ATTRIBUTE_CLOCK_RATE,
                                             current_device));
        checkCudaErrors(cuDeviceComputeCapability(&major, &minor, current_device));

        if (major == 9999 && minor == 9999)
        {
            sm_per_multiproc = 1;
        }
        else
        {
            sm_per_multiproc = _ConvertSMVer2Cores(major, minor);
        }

        int compute_perf  = multiProcessorCount * sm_per_multiproc * clockRate;

        if (compute_perf  > max_compute_perf)
        {
            // If we find GPU with SM major > 2, search only these
            if (best_SM_arch > 2)
            {
                // If our device==dest_SM_arch, choose this, or else pass
                if (major == best_SM_arch)
                {
                    max_compute_perf  = compute_perf;
                    max_perf_device   = current_device;
                }
            }
            else
            {
                max_compute_perf  = compute_perf;
                max_perf_device   = current_device;
            }
        }

        ++current_device;
    }

    return max_perf_device;
}

// General initialization call to pick the best CUDA Device
inline CUdevice findCudaDevice(int argc, char **argv, int *p_devID)
{
    CUdevice cuDevice;
    int devID = 0;

    // If the command-line has a device number specified, use it
    if (checkCmdLineFlag(argc, (const char **)argv, "device"))
    {
        devID = cudaDeviceInit(argc, argv);

        if (devID < 0)
        {
            printf("exiting...\n");
            exit(EXIT_SUCCESS);
        }
    }
    else
    {
        // Otherwise pick the device with highest Gflops/s
        char name[100];
        devID = getMaxGflopsDeviceId();
        checkCudaErrors(cuDeviceGet(&cuDevice, devID));
        cuDeviceGetName(name, 100, cuDevice);
        printf("> Using CUDA Device [%d]: %s\n", devID, name);
    }

    cuDeviceGet(&cuDevice, devID);

    if (p_devID)
    {
        *p_devID = devID;
    }

    return cuDevice;
}
// end of CUDA Helper Functions

// Host code
int main(int argc, char **argv)
{
    pArgc = &argc;
    pArgv = argv;

    printf("Vector Addition (Driver API)\n");
    int N = 50000, devID = 0;
    size_t  size = N * sizeof(float);

    CUresult error;
    ParseArguments(argc, argv);

    // Initialize
    checkCudaErrors(cuInit(0));

    // This assumes that the user is attempting to specify a explicit device -device=n
    if (argc > 1)
    {
        bool bFound = false;

        for (int param=0; param < argc; param++)
        {
            int string_start = 0;

            while (argv[param][string_start] == '-')
            {
                string_start++;
            }

            char *string_argv = &argv[param][string_start];

            if (!strncmp(string_argv, "device", 6))
            {
                int len=(int)strlen(string_argv);

                while (string_argv[len] != '=')
                {
                    len--;
                }

                devID = atoi(&string_argv[++len]);
                bFound = true;
            }

            if (bFound)
            {
                break;
            }
        }
    }

    // Get number of devices supporting CUDA
    int deviceCount = 0;
    error = cuDeviceGetCount(&deviceCount);

    if (error != CUDA_SUCCESS)
    {
        Cleanup(false);
    }

    if (deviceCount == 0)
    {
        printf("There is no device supporting CUDA.\n");
        Cleanup(false);
    }

    if (devID < 0)
    {
        devID = 0;
    }

    if (devID > deviceCount-1)
    {
        fprintf(stderr, "(Device=%d) invalid GPU device.  %d GPU device(s) detected.\nexiting...\n", devID, deviceCount);
        CleanupNoFailure();
        exit(EXIT_SUCCESS);
    }
    else
    {
        int major, minor;
        char deviceName[100];
        checkCudaErrors(cuDeviceComputeCapability(&major, &minor, devID));
        checkCudaErrors(cuDeviceGetName(deviceName, 256, devID));
        printf("> Using Device %d: \"%s\" with Compute %d.%d capability\n", devID, deviceName, major, minor);
    }

    // pick up device with zero ordinal (default, or devID)
    error = cuDeviceGet(&cuDevice, devID);

    if (error != CUDA_SUCCESS)
    {
        Cleanup(false);
    }

    // Create context
    error = cuCtxCreate(&cuContext, 0, cuDevice);

    if (error != CUDA_SUCCESS)
    {
        Cleanup(false);
    }
// CUFFT add
    cudaSetDevice(devID); // should cause CUDA runtime to use the same context
    cufftComplex *cdata;
    cudaMalloc((void **)&cdata, 1024*sizeof(cufftComplex));
    cufftHandle plan;
    cufftResult cstat = cufftPlan1d(&plan, 1024, CUFFT_C2C, 1);
    if (cstat != CUFFT_SUCCESS) printf("CUFFT1 cstat = %d\n", (int)cstat);
    printf("call to cufftPlan1d complete \n");
// 
    // first search for the module path before we load the results
    string module_path, ptx_source;

    if (!findModulePath(PTX_FILE, module_path, argv, ptx_source))
    {
        if (!findModulePath("vectorAdd_kernel.cubin", module_path, argv, ptx_source))
        {
            printf("> findModulePath could not find <vectorAdd> ptx or cubin\n");
            Cleanup(false);
        }
    }
    else
    {
        printf("> initCUDA loading module: <%s>\n", module_path.c_str());
    }

    // Create module from binary file (PTX or CUBIN)
    if (module_path.rfind("ptx") != string::npos)
    {
        // in this branch we use compilation with parameters
        const unsigned int jitNumOptions = 3;
        CUjit_option *jitOptions = new CUjit_option[jitNumOptions];
        void **jitOptVals = new void *[jitNumOptions];

        // set up size of compilation log buffer
        jitOptions[0] = CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES;
        int jitLogBufferSize = 1024;
        jitOptVals[0] = (void *)(size_t)jitLogBufferSize;

        // set up pointer to the compilation log buffer
        jitOptions[1] = CU_JIT_INFO_LOG_BUFFER;
        char *jitLogBuffer = new char[jitLogBufferSize];
        jitOptVals[1] = jitLogBuffer;

        // set up pointer to set the Maximum # of registers for a particular kernel
        jitOptions[2] = CU_JIT_MAX_REGISTERS;
        int jitRegCount = 32;
        jitOptVals[2] = (void *)(size_t)jitRegCount;

        error = cuModuleLoadDataEx(&cuModule, ptx_source.c_str(), jitNumOptions, jitOptions, (void **)jitOptVals);
        printf("call to cuModuleLoadDataEx complete \n");

        printf("> PTX JIT log:\n%s\n", jitLogBuffer);
    }
    else
    {
        error = cuModuleLoad(&cuModule, module_path.c_str());
    }

    if (error != CUDA_SUCCESS)
    {
        Cleanup(false);
    }

    // Get function handle from module
    error = cuModuleGetFunction(&vecAdd_kernel, cuModule, "VecAdd_kernel");

    if (error != CUDA_SUCCESS)
    {
        Cleanup(false);
    }
// CUFFT Add
    cstat = cufftExecC2C(plan, cdata, cdata, CUFFT_FORWARD);
    if (cstat != CUFFT_SUCCESS) printf("CUFFT2 cstat = %d\n", (int)cstat);
    printf("call to cufftExecC2C complete \n");
//
    // Allocate input vectors h_A and h_B in host memory
    h_A = (float *)malloc(size);

    if (h_A == 0)
    {
        Cleanup(false);
    }

    h_B = (float *)malloc(size);

    if (h_B == 0)
    {
        Cleanup(false);
    }

    h_C = (float *)malloc(size);

    if (h_C == 0)
    {
        Cleanup(false);
    }

    // Initialize input vectors
    RandomInit(h_A, N);
    RandomInit(h_B, N);

    // Allocate vectors in device memory
    error = cuMemAlloc(&d_A, size);

    if (error != CUDA_SUCCESS)
    {
        Cleanup(false);
    }

    error = cuMemAlloc(&d_B, size);

    if (error != CUDA_SUCCESS)
    {
        Cleanup(false);
    }

    error = cuMemAlloc(&d_C, size);

    if (error != CUDA_SUCCESS)
    {
        Cleanup(false);
    }

    // Copy vectors from host memory to device memory
    error = cuMemcpyHtoD(d_A, h_A, size);

    if (error != CUDA_SUCCESS)
    {
        Cleanup(false);
    }

    error = cuMemcpyHtoD(d_B, h_B, size);

    if (error != CUDA_SUCCESS)
    {
        Cleanup(false);
    }

#if 1

    if (1)
    {
        // This is the new CUDA 4.0 API for Kernel Parameter Passing and Kernel Launch (simpler method)

        // Grid/Block configuration
        int threadsPerBlock = 256;
        int blocksPerGrid   = (N + threadsPerBlock - 1) / threadsPerBlock;

        void *args[] = { &d_A, &d_B, &d_C, &N };

        // Launch the CUDA kernel
        error = cuLaunchKernel(vecAdd_kernel,  blocksPerGrid, 1, 1,
                               threadsPerBlock, 1, 1,
                               0,
                               NULL, args, NULL);

        if (error != CUDA_SUCCESS)
        {
            Cleanup(false);
        }
    }
    else
    {
        // This is the new CUDA 4.0 API for Kernel Parameter Passing and Kernel Launch (advanced method)
        int offset = 0;
        void *argBuffer[16];
        *((CUdeviceptr *)&argBuffer[offset]) = d_A;
        offset += sizeof(d_A);
        *((CUdeviceptr *)&argBuffer[offset]) = d_B;
        offset += sizeof(d_B);
        *((CUdeviceptr *)&argBuffer[offset]) = d_C;
        offset += sizeof(d_C);
        *((int *)&argBuffer[offset]) = N;
        offset += sizeof(N);

        // Grid/Block configuration
        int threadsPerBlock = 256;
        int blocksPerGrid   = (N + threadsPerBlock - 1) / threadsPerBlock;

        // Launch the CUDA kernel
        error = cuLaunchKernel(vecAdd_kernel,  blocksPerGrid, 1, 1,
                               threadsPerBlock, 1, 1,
                               0,
                               NULL, NULL, argBuffer);

        if (error != CUDA_SUCCESS)
        {
            Cleanup(false);
        }
    }

#else
    {
        char argBuffer[256];

        // pass in launch parameters (not actually de-referencing CUdeviceptr).  CUdeviceptr is
        // storing the value of the parameters
        *((CUdeviceptr *)&argBuffer[offset]) = d_A;
        offset += sizeof(d_A);
        *((CUdeviceptr *)&argBuffer[offset]) = d_B;
        offset += sizeof(d_B);
        *((CUdeviceptr *)&argBuffer[offset]) = d_C;
        offset += sizeof(d_C);
        *((int *)&argBuffer[offset]) = N;
        offset += sizeof(N);

        void *kernel_launch_config[5] =
        {
            CU_LAUNCH_PARAM_BUFFER_POINTER, argBuffer,
            CU_LAUNCH_PARAM_BUFFER_SIZE,    &offset,
            CU_LAUNCH_PARAM_END
        };

        // Grid/Block configuration
        int threadsPerBlock = 256;
        int blocksPerGrid   = (N + threadsPerBlock - 1) / threadsPerBlock;

        // Launch the CUDA kernel
        error = cuLaunchKernel(vecAdd_kernel,  blocksPerGrid, 1, 1,
                               threadsPerBlock, 1, 1,
                               0, 0,
                               NULL, (void **)&kernel_launch_config);

        if (error != CUDA_SUCCESS)
        {
            Cleanup(false);
        }
    }
#endif

#ifdef _DEBUG
    error = cuCtxSynchronize();

    if (error != CUDA_SUCCESS)
    {
        Cleanup(false);
    }

#endif

    // Copy result from device memory to host memory
    // h_C contains the result in host memory
    error = cuMemcpyDtoH(h_C, d_C, size);

    if (error != CUDA_SUCCESS)
    {
        Cleanup(false);
    }

    // Verify result
    int i;

    for (i = 0; i < N; ++i)
    {
        float sum = h_A[i] + h_B[i];

        if (fabs(h_C[i] - sum) > 1e-7f)
        {
            break;
        }
    }

    printf("%s\n", (i==N) ? "Result = PASS" : "Result = FAIL");

    cuProfilerStop();

    exit((i==N) ? EXIT_SUCCESS : EXIT_FAILURE);
}

CUresult CleanupNoFailure()
{
    CUresult error;

    // Free device memory
    if (d_A)
    {
        error = cuMemFree(d_A);
    }

    if (d_B)
    {
        error = cuMemFree(d_B);
    }

    if (d_C)
    {
        error = cuMemFree(d_C);
    }

    // Free host memory
    if (h_A)
    {
        free(h_A);
    }

    if (h_B)
    {
        free(h_B);
    }

    if (h_C)
    {
        free(h_C);
    }

    error = cuCtxDestroy(cuContext);

    return error;
}

void Cleanup(bool noError)
{
    CUresult error;
    error = CleanupNoFailure();

    if (!noError || error != CUDA_SUCCESS)
    {
        printf("Function call failed\nFAILED\n");
        exit(EXIT_FAILURE);
    }

    if (!noprompt)
    {
        printf("\nPress ENTER to exit...\n");
        fflush(stdout);
        fflush(stderr);
        getchar();
    }
}

// Allocates an array with random float entries.
void RandomInit(float *data, int n)
{
    for (int i = 0; i < n; ++i)
    {
        data[i] = rand() / (float)RAND_MAX;
    }
}

bool inline
findModulePath(const char *module_file, string &module_path, char **argv, string &ptx_source)
{
    char *actual_path = sdkFindFilePath(module_file, argv[0]);

    if (actual_path)
    {
        module_path = actual_path;
    }
    else
    {
        printf("> findModulePath file not found: <%s> \n", module_file);
        return false;
    }

    if (module_path.empty())
    {
        printf("> findModulePath could not find file: <%s> \n", module_file);
        return false;
    }
    else
    {
        printf("> findModulePath found file at <%s>\n", module_path.c_str());

        if (module_path.rfind(".ptx") != string::npos)
        {
            FILE *fp = fopen(module_path.c_str(), "rb");
            fseek(fp, 0, SEEK_END);
            int file_size = ftell(fp);
            char *buf = new char[file_size+1];
            fseek(fp, 0, SEEK_SET);
            fread(buf, sizeof(char), file_size, fp);
            fclose(fp);
            buf[file_size] = '

/*

Copyright 1993-2015 NVIDIA Corporation. All rights reserved.
Please refer to the NVIDIA end user license agreement (EULA) associated
with this source code for terms and conditions that govern your use of
this software. Any use, reproduction, disclosure, or distribution of
this software and related documentation outside the terms of the EULA
is strictly prohibited.

*/

/* Vector addition: C = A + B.
*

This sample is a very basic sample that implements element by element
vector addition. It is the same as the sample illustrating Chapter 3
of the programming guide with some additions like error checking.

*/

// Includes
include <stdio.h>
include <string.h>
include
include

// includes, project
include <helper_functions.h>
include <helper_cuda.h>

// includes, CUDA
include <cuda.h>
include <cudaProfiler.h>
include <builtin_types.h>
include <drvapi_error_string.h>
include <cuda_runtime_api.h>
include <cufft.h>

using namespace std;

// Variables
CUdevice cuDevice;
CUcontext cuContext;
CUmodule cuModule;
CUfunction vecAdd_kernel;
float *h_A;
float *h_B;
float *h_C;
CUdeviceptr d_A;
CUdeviceptr d_B;
CUdeviceptr d_C;
bool noprompt = false;

// Functions
void Cleanup(bool);
CUresult CleanupNoFailure();
void RandomInit(float *, int);
bool findModulePath(const char *, string &, char **, string &);
void ParseArguments(int, char **);

int *pArgc = NULL;
char **pArgv = NULL;

//define input ptx file for different platforms
#if defined(_WIN64) || defined(LP64)
define PTX_FILE “vectorAdd_kernel64.ptx”
else
define PTX_FILE “vectorAdd_kernel32.ptx”
endif

////////////////////////////////////////////////////////////////////////////////
// These are CUDA Helper functions

// This will output the proper CUDA error strings in the event that a CUDA host call returns an error
define checkCudaErrors(err) __checkCudaErrors (err, FILE, LINE)

// These are the inline versions for all of the SDK helper functions
inline void __checkCudaErrors(CUresult err, const char *file, const int line)
{
if (CUDA_SUCCESS != err)
{
fprintf(stderr, “checkCudaErrors() Driver API error = %04d "%s" from file <%s>, line %i.\n”,
err, getCudaDrvErrorString(err), file, line);
exit(EXIT_FAILURE);
}
}

inline int cudaDeviceInit(int ARGC, char **ARGV)
{
int cuDevice = 0;
int deviceCount = 0;
CUresult err = cuInit(0);

if (CUDA_SUCCESS == err)
{
    checkCudaErrors(cuDeviceGetCount(&deviceCount));
}

if (deviceCount == 0)
{
    fprintf(stderr, "cudaDeviceInit error: no devices supporting CUDA\n");
    exit(EXIT_FAILURE);
}

int dev = 0;
dev = getCmdLineArgumentInt(ARGC, (const char **) ARGV, "device=");

if (dev < 0)
{
    dev = 0;
}

if (dev > deviceCount-1)
{
    fprintf(stderr, "\n");
    fprintf(stderr, ">> %d CUDA capable GPU device(s) detected. <<\n", deviceCount);
    fprintf(stderr, ">> cudaDeviceInit (-device=%d) is not a valid GPU device. <<\n", dev);
    fprintf(stderr, "\n");
    return -dev;
}

checkCudaErrors(cuDeviceGet(&cuDevice, dev));
char name[100];
cuDeviceGetName(name, 100, cuDevice);

if (checkCmdLineFlag(ARGC, (const char **) ARGV, "quiet") == false)
{
    printf("> Using CUDA Device [%d]: %s\n", dev, name);
}

return dev;

}

// This function returns the best GPU based on performance
inline int getMaxGflopsDeviceId()
{
CUdevice current_device = 0, max_perf_device = 0;
int device_count = 0, sm_per_multiproc = 0;
int max_compute_perf = 0, best_SM_arch = 0;
int major = 0, minor = 0, multiProcessorCount, clockRate;

cuInit(0);
checkCudaErrors(cuDeviceGetCount(&device_count));

// Find the best major SM Architecture GPU device
while (current_device < device_count)
{
    checkCudaErrors(cuDeviceComputeCapability(&major, &minor, current_device));

    if (major > 0 && major < 9999)
    {
        best_SM_arch = MAX(best_SM_arch, major);
    }

    current_device++;
}

// Find the best CUDA capable GPU device
current_device = 0;

while (current_device < device_count)
{
    checkCudaErrors(cuDeviceGetAttribute(&multiProcessorCount,
                                         CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT,
                                         current_device));
    checkCudaErrors(cuDeviceGetAttribute(&clockRate,
                                         CU_DEVICE_ATTRIBUTE_CLOCK_RATE,
                                         current_device));
    checkCudaErrors(cuDeviceComputeCapability(&major, &minor, current_device));

    if (major == 9999 && minor == 9999)
    {
        sm_per_multiproc = 1;
    }
    else
    {
        sm_per_multiproc = _ConvertSMVer2Cores(major, minor);
    }

    int compute_perf  = multiProcessorCount * sm_per_multiproc * clockRate;

    if (compute_perf  > max_compute_perf)
    {
        // If we find GPU with SM major > 2, search only these
        if (best_SM_arch > 2)
        {
            // If our device==dest_SM_arch, choose this, or else pass
            if (major == best_SM_arch)
            {
                max_compute_perf  = compute_perf;
                max_perf_device   = current_device;
            }
        }
        else
        {
            max_compute_perf  = compute_perf;
            max_perf_device   = current_device;
        }
    }

    ++current_device;
}

return max_perf_device;

}

// General initialization call to pick the best CUDA Device
inline CUdevice findCudaDevice(int argc, char **argv, int *p_devID)
{
CUdevice cuDevice;
int devID = 0;

// If the command-line has a device number specified, use it
if (checkCmdLineFlag(argc, (const char **)argv, "device"))
{
    devID = cudaDeviceInit(argc, argv);

    if (devID < 0)
    {
        printf("exiting...\n");
        exit(EXIT_SUCCESS);
    }
}
else
{
    // Otherwise pick the device with highest Gflops/s
    char name[100];
    devID = getMaxGflopsDeviceId();
    checkCudaErrors(cuDeviceGet(&cuDevice, devID));
    cuDeviceGetName(name, 100, cuDevice);
    printf("> Using CUDA Device [%d]: %s\n", devID, name);
}

cuDeviceGet(&cuDevice, devID);

if (p_devID)
{
    *p_devID = devID;
}

return cuDevice;

}
// end of CUDA Helper Functions

// Host code
int main(int argc, char **argv)
{
pArgc = &argc;
pArgv = argv;

printf("Vector Addition (Driver API)\n");
int N = 50000, devID = 0;
size_t  size = N * sizeof(float);

CUresult error;
ParseArguments(argc, argv);

// Initialize
checkCudaErrors(cuInit(0));

// This assumes that the user is attempting to specify a explicit device -device=n
if (argc > 1)
{
    bool bFound = false;

    for (int param=0; param < argc; param++)
    {
        int string_start = 0;

        while (argv[param][string_start] == '-')
        {
            string_start++;
        }

        char *string_argv = &argv[param][string_start];

        if (!strncmp(string_argv, "device", 6))
        {
            int len=(int)strlen(string_argv);

            while (string_argv[len] != '=')
            {
                len--;
            }

            devID = atoi(&string_argv[++len]);
            bFound = true;
        }

        if (bFound)
        {
            break;
        }
    }
}

// Get number of devices supporting CUDA
int deviceCount = 0;
error = cuDeviceGetCount(&deviceCount);

if (error != CUDA_SUCCESS)
{
    Cleanup(false);
}

if (deviceCount == 0)
{
    printf("There is no device supporting CUDA.\n");
    Cleanup(false);
}

if (devID < 0)
{
    devID = 0;
}

if (devID > deviceCount-1)
{
    fprintf(stderr, "(Device=%d) invalid GPU device.  %d GPU device(s) detected.\nexiting...\n", devID, deviceCount);
    CleanupNoFailure();
    exit(EXIT_SUCCESS);
}
else
{
    int major, minor;
    char deviceName[100];
    checkCudaErrors(cuDeviceComputeCapability(&major, &minor, devID));
    checkCudaErrors(cuDeviceGetName(deviceName, 256, devID));
    printf("> Using Device %d: \"%s\" with Compute %d.%d capability\n", devID, deviceName, major, minor);
}

// pick up device with zero ordinal (default, or devID)
error = cuDeviceGet(&cuDevice, devID);

if (error != CUDA_SUCCESS)
{
    Cleanup(false);
}

// Create context
error = cuCtxCreate(&cuContext, 0, cuDevice);

if (error != CUDA_SUCCESS)
{
    Cleanup(false);
}

// CUFFT add
cudaSetDevice(devID); // should cause CUDA runtime to use the same context
cufftComplex cdata;
cudaMalloc((void **)&cdata, 1024sizeof(cufftComplex));
cufftHandle plan;
cufftResult cstat = cufftPlan1d(&plan, 1024, CUFFT_C2C, 1);
if (cstat != CUFFT_SUCCESS) printf(“CUFFT1 cstat = %d\n”, (int)cstat);
printf(“call to cufftPlan1d complete \n”);
//
// first search for the module path before we load the results
string module_path, ptx_source;

if (!findModulePath(PTX_FILE, module_path, argv, ptx_source))
{
    if (!findModulePath("vectorAdd_kernel.cubin", module_path, argv, ptx_source))
    {
        printf("> findModulePath could not find <vectorAdd> ptx or cubin\n");
        Cleanup(false);
    }
}
else
{
    printf("> initCUDA loading module: <%s>\n", module_path.c_str());
}

// Create module from binary file (PTX or CUBIN)
if (module_path.rfind("ptx") != string::npos)
{
    // in this branch we use compilation with parameters
    const unsigned int jitNumOptions = 3;
    CUjit_option *jitOptions = new CUjit_option[jitNumOptions];
    void **jitOptVals = new void *[jitNumOptions];

    // set up size of compilation log buffer
    jitOptions[0] = CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES;
    int jitLogBufferSize = 1024;
    jitOptVals[0] = (void *)(size_t)jitLogBufferSize;

    // set up pointer to the compilation log buffer
    jitOptions[1] = CU_JIT_INFO_LOG_BUFFER;
    char *jitLogBuffer = new char[jitLogBufferSize];
    jitOptVals[1] = jitLogBuffer;

    // set up pointer to set the Maximum # of registers for a particular kernel
    jitOptions[2] = CU_JIT_MAX_REGISTERS;
    int jitRegCount = 32;
    jitOptVals[2] = (void *)(size_t)jitRegCount;

    error = cuModuleLoadDataEx(&cuModule, ptx_source.c_str(), jitNumOptions, jitOptions, (void **)jitOptVals);
    printf("call to cuModuleLoadDataEx complete \n");

    printf("> PTX JIT log:\n%s\n", jitLogBuffer);
}
else
{
    error = cuModuleLoad(&cuModule, module_path.c_str());
}

if (error != CUDA_SUCCESS)
{
    Cleanup(false);
}

// Get function handle from module
error = cuModuleGetFunction(&vecAdd_kernel, cuModule, "VecAdd_kernel");

if (error != CUDA_SUCCESS)
{
    Cleanup(false);
}

// CUFFT Add
cstat = cufftExecC2C(plan, cdata, cdata, CUFFT_FORWARD);
if (cstat != CUFFT_SUCCESS) printf(“CUFFT2 cstat = %d\n”, (int)cstat);
printf(“call to cufftExecC2C complete \n”);
//
// Allocate input vectors h_A and h_B in host memory
h_A = (float *)malloc(size);

if (h_A == 0)
{
    Cleanup(false);
}

h_B = (float *)malloc(size);

if (h_B == 0)
{
    Cleanup(false);
}

h_C = (float *)malloc(size);

if (h_C == 0)
{
    Cleanup(false);
}

// Initialize input vectors
RandomInit(h_A, N);
RandomInit(h_B, N);

// Allocate vectors in device memory
error = cuMemAlloc(&d_A, size);

if (error != CUDA_SUCCESS)
{
    Cleanup(false);
}

error = cuMemAlloc(&d_B, size);

if (error != CUDA_SUCCESS)
{
    Cleanup(false);
}

error = cuMemAlloc(&d_C, size);

if (error != CUDA_SUCCESS)
{
    Cleanup(false);
}

// Copy vectors from host memory to device memory
error = cuMemcpyHtoD(d_A, h_A, size);

if (error != CUDA_SUCCESS)
{
    Cleanup(false);
}

error = cuMemcpyHtoD(d_B, h_B, size);

if (error != CUDA_SUCCESS)
{
    Cleanup(false);
}

#if 1

if (1)
{
    // This is the new CUDA 4.0 API for Kernel Parameter Passing and Kernel Launch (simpler method)

    // Grid/Block configuration
    int threadsPerBlock = 256;
    int blocksPerGrid   = (N + threadsPerBlock - 1) / threadsPerBlock;

    void *args[] = { &d_A, &d_B, &d_C, &N };

    // Launch the CUDA kernel
    error = cuLaunchKernel(vecAdd_kernel,  blocksPerGrid, 1, 1,
                           threadsPerBlock, 1, 1,
                           0,
                           NULL, args, NULL);

    if (error != CUDA_SUCCESS)
    {
        Cleanup(false);
    }
}
else
{
    // This is the new CUDA 4.0 API for Kernel Parameter Passing and Kernel Launch (advanced method)
    int offset = 0;
    void *argBuffer[16];
    *((CUdeviceptr *)&argBuffer[offset]) = d_A;
    offset += sizeof(d_A);
    *((CUdeviceptr *)&argBuffer[offset]) = d_B;
    offset += sizeof(d_B);
    *((CUdeviceptr *)&argBuffer[offset]) = d_C;
    offset += sizeof(d_C);
    *((int *)&argBuffer[offset]) = N;
    offset += sizeof(N);

    // Grid/Block configuration
    int threadsPerBlock = 256;
    int blocksPerGrid   = (N + threadsPerBlock - 1) / threadsPerBlock;

    // Launch the CUDA kernel
    error = cuLaunchKernel(vecAdd_kernel,  blocksPerGrid, 1, 1,
                           threadsPerBlock, 1, 1,
                           0,
                           NULL, NULL, argBuffer);

    if (error != CUDA_SUCCESS)
    {
        Cleanup(false);
    }
}

else
{
char argBuffer[256];

    // pass in launch parameters (not actually de-referencing CUdeviceptr).  CUdeviceptr is
    // storing the value of the parameters
    *((CUdeviceptr *)&argBuffer[offset]) = d_A;
    offset += sizeof(d_A);
    *((CUdeviceptr *)&argBuffer[offset]) = d_B;
    offset += sizeof(d_B);
    *((CUdeviceptr *)&argBuffer[offset]) = d_C;
    offset += sizeof(d_C);
    *((int *)&argBuffer[offset]) = N;
    offset += sizeof(N);

    void *kernel_launch_config[5] =
    {
        CU_LAUNCH_PARAM_BUFFER_POINTER, argBuffer,
        CU_LAUNCH_PARAM_BUFFER_SIZE,    &offset,
        CU_LAUNCH_PARAM_END
    };

    // Grid/Block configuration
    int threadsPerBlock = 256;
    int blocksPerGrid   = (N + threadsPerBlock - 1) / threadsPerBlock;

    // Launch the CUDA kernel
    error = cuLaunchKernel(vecAdd_kernel,  blocksPerGrid, 1, 1,
                           threadsPerBlock, 1, 1,
                           0, 0,
                           NULL, (void **)&kernel_launch_config);

    if (error != CUDA_SUCCESS)
    {
        Cleanup(false);
    }
}

endif

ifdef _DEBUG
error = cuCtxSynchronize();

if (error != CUDA_SUCCESS)
{
    Cleanup(false);
}

endif

// Copy result from device memory to host memory
// h_C contains the result in host memory
error = cuMemcpyDtoH(h_C, d_C, size);

if (error != CUDA_SUCCESS)
{
    Cleanup(false);
}

// Verify result
int i;

for (i = 0; i < N; ++i)
{
    float sum = h_A[i] + h_B[i];

    if (fabs(h_C[i] - sum) > 1e-7f)
    {
        break;
    }
}

printf("%s\n", (i==N) ? "Result = PASS" : "Result = FAIL");

cuProfilerStop();

exit((i==N) ? EXIT_SUCCESS : EXIT_FAILURE);

}

CUresult CleanupNoFailure()
{
CUresult error;

// Free device memory
if (d_A)
{
    error = cuMemFree(d_A);
}

if (d_B)
{
    error = cuMemFree(d_B);
}

if (d_C)
{
    error = cuMemFree(d_C);
}

// Free host memory
if (h_A)
{
    free(h_A);
}

if (h_B)
{
    free(h_B);
}

if (h_C)
{
    free(h_C);
}

error = cuCtxDestroy(cuContext);

return error;

}

void Cleanup(bool noError)
{
CUresult error;
error = CleanupNoFailure();

if (!noError || error != CUDA_SUCCESS)
{
    printf("Function call failed\nFAILED\n");
    exit(EXIT_FAILURE);
}

if (!noprompt)
{
    printf("\nPress ENTER to exit...\n");
    fflush(stdout);
    fflush(stderr);
    getchar();
}

}

// Allocates an array with random float entries.
void RandomInit(float *data, int n)
{
for (int i = 0; i < n; ++i)
{
data[i] = rand() / (float)RAND_MAX;
}
}

bool inline
findModulePath(const char *module_file, string &module_path, char **argv, string &ptx_source)
{
char *actual_path = sdkFindFilePath(module_file, argv[0]);

if (actual_path)
{
    module_path = actual_path;
}
else
{
    printf("> findModulePath file not found: <%s> \n", module_file);
    return false;
}

if (module_path.empty())
{
    printf("> findModulePath could not find file: <%s> \n", module_file);
    return false;
}
else
{
    printf("> findModulePath found file at <%s>\n", module_path.c_str());

    if (module_path.rfind(".ptx") != string::npos)
    {
        FILE *fp = fopen(module_path.c_str(), "rb");
        fseek(fp, 0, SEEK_END);
        int file_size = ftell(fp);
        char *buf = new char[file_size+1];
        fseek(fp, 0, SEEK_SET);
        fread(buf, sizeof(char), file_size, fp);
        fclose(fp);
        buf[file_size] = '\0';
        ptx_source = buf;
        delete[] buf;
    }

    return true;
}

}

// Parse program arguments
void ParseArguments(int argc, char **argv)
{
for (int i = 0; i < argc; ++i)
{
if (strcmp(argv[i], “–noprompt”) == 0 ||
strcmp(argv[i], “-noprompt”) == 0)
{
noprompt = true;
break;
}
}
}

';
            ptx_source = buf;
            delete[] buf;
        }

        return true;
    }
}

// Parse program arguments
void ParseArguments(int argc, char **argv)
{
    for (int i = 0; i < argc; ++i)
    {
        if (strcmp(argv[i], "--noprompt") == 0 ||
            strcmp(argv[i], "-noprompt") == 0)
        {
            noprompt = true;
            break;
        }
    }
}

(no change to the ptx kernel)

compile/run:

$ nvcc -ptx vectorAdd_kernel.cu -o vectorAdd_kernel64.ptx
$ nvcc -I/usr/local/cuda/samples/common/inc -o vectorAddDrv vectorAddDrv.cpp -lcufft -lcuda
$ cuda-memcheck ./vectorAddDrv
========= CUDA-MEMCHECK
Vector Addition (Driver API)
> Using Device 0: "Quadro 5000" with Compute 2.0 capability
call to cufftPlan1d complete
> findModulePath found file at <./vectorAdd_kernel64.ptx>
> initCUDA loading module: <./vectorAdd_kernel64.ptx>
call to cuModuleLoadDataEx complete
> PTX JIT log:

call to cufftExecC2C complete
Result = PASS
========= ERROR SUMMARY: 0 errors
$

Topic		Replies	Views
Invalid __global__ write of size 4 Error on kernel Launch CUDA Programming and Performance	7	6115	March 18, 2013
Help!! I can't get my NVidia GeForce GT 525M to load in a single CUDA PTX kernel!! CUDA Programming and Performance	11	5793	November 16, 2012
CUFFT_INVALID_DEVICE on cufftPlan1d in NVIDIA's Simple CUFFT example GPU-Accelerated Libraries	6	3798	December 15, 2014
cufftXt batch 1D GPU-Accelerated Libraries	12	2149	October 15, 2019
cuFFT and fftw CUDA Programming and Performance	10	4145	August 25, 2010
[SOLVED] cuFFT data storage, maybe I'm operating on the wrong elements CUDA Programming and Performance	16	2114	June 11, 2022
Fourth FFT in cuda throws illegal memory access GPU-Accelerated Libraries cuda	3	553	December 10, 2020
Using Thrust to sort Unified Memory Buffer? GPU-Accelerated Libraries	8	5070	May 7, 2015
New Features in CUDA 7.5 Technical Blog	66	1066	August 10, 2016
CUFFT_INTERNAL_ERROR during creation of a 1D Plan in CUFFT GPU-Accelerated Libraries cuda , cufft	11	3728	October 19, 2022

CUFFT_INVALID_PLAN after loading other ptx code with cuModuleLoadDataEx

Related topics