test_ptxas_verbose.c

test_ptxas_verbose.c

#if __APPLE__
#include "CUDA/cuda.h"
#else
#include "cuda.h"
#endif

const char ptx[] = "//\n// Generated by NVIDIA NVVM Compiler\n// Compiler built on Thu May 16 17:15:21 2013 (1368742521)\n// Cuda compilation tools, release 5.5, V5.5.0\n//\n\n.version 3.2\n.target sm_20\n.address_size 64\n\n\n.visible .entry sum(\n\t.param .u64 sum_param_0\n)\n{\n\t.reg .s32 \t%r<3>;\n\t.reg .s64 \t%rd<5>;\n\n\n\tld.param.u64 \t%rd1, [sum_param_0];\n\tcvta.to.global.u64 \t%rd2, %rd1;\n\tld.global.u64 \t%rd3, [%rd2];\n\tcvta.to.global.u64 \t%rd4, %rd3;\n\tld.global.u32 \t%r1, [%rd4];\n\tadd.s32 \t%r2, %r1, 1;\n\tst.global.u32 \t[%rd4], %r2;\n\tret;\n}\n\n\n\x00";

int main(int argc, char **argv){
    CUdevice device;
    CUcontext context;
    CUmodule cudaModule;

    cuInit(0);
    cuDeviceGet(&device, 0);

    puts("device inited");
    cuCtxCreate(&context, 0, device);
    puts("context inited");


    char linkerInfo[1024];
    char linkerErrors[1024];

    linkerInfo[0] = 0;
    linkerErrors[0] = 0;

    CUjit_option linkerOptions[] = {
    CU_JIT_INFO_LOG_BUFFER,
    CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES,
    CU_JIT_ERROR_LOG_BUFFER,
    CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES,
    CU_JIT_LOG_VERBOSE
    };

    void *linkerOptionValues[] = {
        linkerInfo,
        (void*)(1024),
        linkerErrors,
        (void*)(1024),
        (void*)(1)
    };

    puts("load module");
    cuModuleLoadDataEx(&cudaModule, ptx,
                       sizeof(linkerOptions)/sizeof(CUjit_option),
                       linkerOptions,
                       linkerOptionValues);

    puts(linkerInfo);
    puts(linkerErrors);
    puts("cleanup");
    cuModuleUnload(cudaModule);
    cuCtxDestroy(context);
    
    puts("end");
    return 0;
}