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");
  }
}
Exemplo n.º 2
0
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;
					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);
}