String c_AsyncFunctionWaitHandle::getName() {
switch (getState()) {
case STATE_BLOCKED:
case STATE_READY:
case STATE_RUNNING: {
auto func = actRec()->func();
if (!func->cls() ||
func->cls()->attrs() & AttrNoOverride ||
actRec()->localsDecRefd()) {
auto name = func->fullName();
if (func->isClosureBody()) {
const char* p = strchr(name->data(), ':');
if (p) {
return
concat(String(name->data(), p + 1 - name->data(), CopyString),
s__closure_);
} else {
return s__closure_;
}
}
return String{const_cast<StringData*>(name)};
}
auto const cls = actRec()->hasThis() ?
actRec()->getThis()->getVMClass() :
actRec()->getClass();
if (cls == func->cls() && !func->isClosureBody()) {
return String{const_cast<StringData*>(func->fullName())};
}
StrNR funcName(func->isClosureBody() ? s__closure_.get() : func->name());
return concat3(cls->nameStr(), "::", funcName);
}
default:
raise_fatal_error(
"Invariant violation: encountered unexpected state");
}
}
Variant invoke(const char *function, const Variant& params, strhash_t hash /* = -1 */,
bool tryInterp /* = true */, bool fatal /* = true */) {
String funcName(function, CopyString);
return invoke(funcName, params, hash, tryInterp, fatal);
}
int main( int argc, char **argv )
{
CUdevice main_device = 0;
CUcontext main_context = nullptr;
CUmodule mod_vectorAdd = nullptr;
CUfunction fun_vectorAdd = nullptr;
std::string path_vectorAdd( "D:/devel/vectoradd-cuda-driverAPI/vectorAdd.cu" );
std::string ptx_vectorAdd ( "D:/devel/vectoradd-cuda-driverAPI/vectorAdd.ptx" );
CUdeviceptr input_data = 0u;
CUdeviceptr output_data = 0u;
std::size_t problem_size = 1024;
try
{
//Initialize the driver API
check_error( cuInit( 0u ) );
{
int device_count = 0u;
check_error( cuDeviceGetCount( &device_count ) );
if( ! device_count )
{
std::cerr << "No CUDA devices available" << std::endl;
throw CUDA_ERROR_NO_DEVICE;
}
}
check_error( cuDeviceGet( &main_device, 0 ) );
check_error( cuCtxCreate( &main_context, 0, main_device ) );
//Try to manually compile the source file
{
std::stringstream build_command;
build_command <<
"nvcc "
"-ptx "
"-o " << ptx_vectorAdd << " " <<
path_vectorAdd;
if( int build_status = system( build_command.str( ).c_str( ) ) )
{
std::cerr << "Failed to compile source cuda file into a ptx assembly" << std::endl;
throw CUDA_ERROR_UNKNOWN;
}
//Find module entry with assembly
std::string str_assembly;
{
std::ifstream fassembly( ptx_vectorAdd );
if( !fassembly.is_open( ) )
{
std::cerr << "'Vector Add' assembly unavailable" << std::endl;
throw CUDA_ERROR_FILE_NOT_FOUND;
}
fassembly.seekg (0, std::ios::end);
str_assembly.resize( std::string::size_type( fassembly.tellg() ) );
fassembly.seekg (0, std::ios::beg);
fassembly.read( &str_assembly[0], str_assembly.size( ) );
fassembly.close( );
}
auto entry_pos = str_assembly.find( ".entry" );
if( entry_pos == std::string::npos )
{
std::cerr << "No entry point in 'Vector Add'" << std::endl;
throw CUDA_ERROR_INVALID_SOURCE;
}
entry_pos += 6u; //".entry".size( )
auto search_limit = str_assembly.find_first_of( " (", entry_pos );
if( search_limit == std::string::npos )
{
std::cerr << "No entry point in 'Vector Add'" << std::endl;
throw CUDA_ERROR_INVALID_SOURCE;
}
std::string funcName( str_assembly.substr( entry_pos, search_limit ) );
check_error( cuModuleLoad ( &mod_vectorAdd, ptx_vectorAdd.c_str( ) ) );
check_error( cuModuleGetFunction ( &fun_vectorAdd, mod_vectorAdd, funcName.c_str( ) ) );
}
//Play with buffer
cuMemAlloc( &input_data, problem_size * sizeof( float ) );
cuMemAlloc( &output_data, problem_size * sizeof( float ) );
{
int threadsPerBlock = 256;
int blocksPerGrid = (problem_size + threadsPerBlock - 1) / threadsPerBlock;
void* args[] = { &input_data, &output_data, &problem_size };
cuLaunchKernel(
fun_vectorAdd,
blocksPerGrid, 1, 1,
threadsPerBlock, 1, 1,
0,
0,
args,
nullptr);
}
float* result = new float[problem_size];
cuMemcpyDtoH( result, output_data, problem_size * sizeof( float ) );
std::copy(
result, result + problem_size,
std::ostream_iterator<float>(std::cout, ", ") );
delete[] result;
if( output_data ) cuMemFree ( output_data );
if( input_data ) cuMemFree ( input_data );
if( mod_vectorAdd ) cuModuleUnload ( mod_vectorAdd );
if( main_context ) cuCtxDestroy ( main_context );
}
catch( int return_code )
{
if( output_data ) cuMemFree ( output_data );
if( input_data ) cuMemFree ( input_data );
if( mod_vectorAdd ) cuModuleUnload ( mod_vectorAdd );
if( main_context ) cuCtxDestroy ( main_context );
system("PAUSE");
return return_code;
//==============================================================================
void MaterialProgramCreator::parseOperationTag(
const XmlElement& operationTag, GLenum glshader, GLbitfield glshaderbit,
MPString& out)
{
static const char OUT[] = {"out"};
// <id></id>
I id = operationTag.getChildElement("id").getInt();
// <returnType></returnType>
XmlElement retTypeEl = operationTag.getChildElement("returnType");
MPString retType(retTypeEl.getText(), m_alloc);
MPString operationOut(m_alloc);
if(retType != "void")
{
MPString tmp(MPString::toString(id, m_alloc));
operationOut = ANKI_STRL(OUT) + tmp;
}
// <function>functionName</function>
MPString funcName(
operationTag.getChildElement("function").getText(), m_alloc);
// <arguments></arguments>
XmlElement argsEl = operationTag.getChildElementOptional("arguments");
MPStringList argsList(m_alloc);
if(argsEl)
{
// Get all arguments
XmlElement argEl = argsEl.getChildElement("argument");
do
{
MPString arg(argEl.getText(), m_alloc);
// Search for all the inputs and mark the appropriate
Input* input = nullptr;
for(Input& in : m_inputs)
{
// Check that the first part of the string is equal to the
// variable and the following char is '['
if(in.m_name == arg)
{
input = ∈
in.m_shaderReferencedMask = glshaderbit;
break;
}
}
// The argument should be an input variable or an outXX
if(!(input != nullptr
|| std::strncmp(&arg[0], OUT, sizeof(OUT) - 1) == 0))
{
throw ANKI_EXCEPTION("Incorrect argument: %s", &arg[0]);
}
// Add to a list and do something special if instanced
if(input && input->m_instanced)
{
if(glshader == GL_VERTEX_SHADER)
{
argsList.push_back(ANKI_STRL(argEl.getText())
+ "[gl_InstanceID]");
m_instanceIdMask |= glshaderbit;
}
else if(glshader == GL_TESS_CONTROL_SHADER)
{
argsList.push_back(ANKI_STRL(argEl.getText())
+ "[vInstanceId[0]]");
m_instanceIdMask |= glshaderbit;
}
else if(glshader == GL_TESS_EVALUATION_SHADER)
{
argsList.push_back(ANKI_STRL(argEl.getText())
+ "[commonPatch.instanceId]");
m_instanceIdMask |= glshaderbit;
}
else if(glshader == GL_FRAGMENT_SHADER)
{
argsList.push_back(ANKI_STRL(argEl.getText())
+ "[vInstanceId]");
m_instanceIdMask |= glshaderbit;
}
else
{
throw ANKI_EXCEPTION(
"Cannot access the instance ID in all shaders");
}
}
else
{
argsList.push_back(MPString(argEl.getText(), m_alloc));
}
// Advance
argEl = argEl.getNextSiblingElement("argument");
} while(argEl);
}
// Now write everything
MPString lines(m_alloc);
lines.reserve(256);
lines += "#if defined(" + funcName + "_DEFINED)";
// Write the defines for the operationOuts
for(const MPString& arg : argsList)
{
if(arg.find(OUT) == 0)
{
lines += " && defined(" + arg + "_DEFINED)";
}
}
lines += "\n";
if(retType != "void")
{
lines += "#\tdefine " + operationOut + "_DEFINED\n\t"
+ retTypeEl.getText() + " " + operationOut + " = ";
}
else
{
lines += "\t";
}
// write the blah = func(args...)
lines += funcName + "(";
lines += argsList.join(", ");
lines += ");\n";
lines += "#endif";
// Done
out = std::move(lines);
}
